revert inadvertent change to .pre-commit-config.yaml

Signed-off-by: Nick Hill <nickhill123@gmail.com>
[Doc] register_kv_caches: views are authoritative, not storage nbytes
2026-07-22 16:12:01 +01:00 · 2026-07-22 15:59:16 +01:00 · 2026-07-22 15:59:16 +01:00 · 2026-07-22 14:40:35 +01:00 · 2026-07-22 08:30:35 +00:00 · 2026-07-22 08:22:15 +00:00
440 changed files with 32474 additions and 5227 deletions
@@ -18,6 +18,8 @@ steps:
  - tests/kernels/quantization/test_cpu_fp8_scaled_mm.py
  - tests/kernels/mamba/cpu/test_cpu_gdn_ops.py
  - tests/kernels/mamba/test_cpu_short_conv.py
+  - tests/kernels/mamba/test_causal_conv1d.py
+  - tests/kernels/mamba/test_mamba_ssm.py
  commands:
    - |
      bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 30m "
@@ -28,7 +30,9 @@ steps:
      pytest -x -v -s tests/kernels/test_onednn.py
      pytest -x -v -s tests/kernels/test_awq_int4_to_int8.py
      pytest -x -v -s tests/kernels/quantization/test_cpu_fp8_scaled_mm.py
-      pytest -x -v -s tests/kernels/mamba/cpu/test_cpu_gdn_ops.py"
+      pytest -x -v -s tests/kernels/mamba/cpu/test_cpu_gdn_ops.py
+      pytest -x -v -s tests/kernels/mamba/test_causal_conv1d.py
+      pytest -x -v -s tests/kernels/mamba/test_mamba_ssm.py"

 # Note: SDE can't be downloaded from CI host because of AWS WAF
 # - label: CPU-Compatibility Tests
@@ -31,8 +31,46 @@ steps:
      - text: "What is the release version?"
        key: release-version

-  - group: "Build Python wheels"
+  - group: "Build CUDA 13.0 Python wheels"
    key: "build-wheels"
+    steps:
+      - label: "Build wheel - aarch64 - CUDA 13.0"
+        depends_on: ~
+        id: build-wheel-arm64-cuda-13-0
+        agents:
+          queue: arm64_cpu_queue_release
+        commands:
+          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=13.0.2 --build-arg torch_cuda_arch_list=\"${CUDA_ARCH_AARCH64}\" --build-arg BUILD_OS=manylinux --build-arg BUILD_BASE_IMAGE=pytorch/manylinuxaarch64-builder:cuda13.0 --tag vllm-ci:build-image --target build --progress plain -f docker/Dockerfile ."
+          - "mkdir artifacts"
+          - "docker run --rm -v $(pwd)/artifacts:/artifacts_host vllm-ci:build-image bash -c 'cp -r dist /artifacts_host && chmod -R a+rw /artifacts_host'"
+          - "bash .buildkite/scripts/upload-nightly-wheels.sh"
+          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "s3://vllm-wheels/$$BUILDKITE_COMMIT/$(cd artifacts/dist && echo *.whl)" release-wheels'
+        env:
+          DOCKER_BUILDKIT: "1"
+
+      - label: "Build wheel - x86_64 - CUDA 13.0"
+        depends_on: ~
+        id: build-wheel-x86-cuda-13-0
+        agents:
+          queue: cpu_queue_release
+        commands:
+          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=13.0.2 --build-arg torch_cuda_arch_list=\"${CUDA_ARCH_X86}\" --build-arg BUILD_OS=manylinux --build-arg BUILD_BASE_IMAGE=pytorch/manylinux2_28-builder:cuda13.0 --tag vllm-ci:build-image --target build --progress plain -f docker/Dockerfile ."
+          - "mkdir artifacts"
+          - "docker run --rm -v $(pwd)/artifacts:/artifacts_host vllm-ci:build-image bash -c 'cp -r dist /artifacts_host && chmod -R a+rw /artifacts_host'"
+          - "bash .buildkite/scripts/upload-nightly-wheels.sh"
+          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "s3://vllm-wheels/$$BUILDKITE_COMMIT/$(cd artifacts/dist && echo *.whl)" release-wheels'
+        env:
+          DOCKER_BUILDKIT: "1"
+
+  - block: "Unblock to build additional Python wheels"
+    depends_on: ~
+    key: block-build-additional-wheels
+    if: build.env("NIGHTLY") != "1"
+
+  - group: "Build additional Python wheels"
+    key: "build-additional-wheels"
+    depends_on: block-build-additional-wheels
+    allow_dependency_failure: true
    steps:
      - label: "Build wheel - aarch64 - CUDA 12.9"
        depends_on: ~
@@ -48,20 +86,6 @@ steps:
        env:
          DOCKER_BUILDKIT: "1"

-      - label: "Build wheel - aarch64 - CUDA 13.0"
-        depends_on: ~
-        id: build-wheel-arm64-cuda-13-0
-        agents:
-          queue: arm64_cpu_queue_release
-        commands:
-          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=13.0.2 --build-arg torch_cuda_arch_list=\"${CUDA_ARCH_AARCH64}\" --build-arg BUILD_OS=manylinux --build-arg BUILD_BASE_IMAGE=pytorch/manylinuxaarch64-builder:cuda13.0 --tag vllm-ci:build-image --target build --progress plain -f docker/Dockerfile ."
-          - "mkdir artifacts"
-          - "docker run --rm -v $(pwd)/artifacts:/artifacts_host vllm-ci:build-image bash -c 'cp -r dist /artifacts_host && chmod -R a+rw /artifacts_host'"
-          - "bash .buildkite/scripts/upload-nightly-wheels.sh"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "s3://vllm-wheels/$$BUILDKITE_COMMIT/$(cd artifacts/dist && echo *.whl)" release-wheels'
-        env:
-          DOCKER_BUILDKIT: "1"
-
      - label: "Build wheel - aarch64 - CPU"
        depends_on: ~
        id: build-wheel-arm64-cpu
@@ -113,7 +137,7 @@ steps:
          - 'mv artifacts/reassembled/wheel "artifacts/dist/$$wheel_name"'
          - "aws sts get-caller-identity"
          - "VLLM_WHEEL_PLATFORM=macos bash .buildkite/scripts/upload-nightly-wheels.sh"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "s3://vllm-wheels/$$BUILDKITE_COMMIT/$(cd artifacts/dist && echo *.whl)"'
+          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "s3://vllm-wheels/$$BUILDKITE_COMMIT/$(cd artifacts/dist && echo *.whl)" release-wheels'
        plugins:
          - aws-assume-role-with-web-identity#v1.6.0:
              role-arn: arn:aws:iam::936637512419:role/vllm-release-macos-wheel-uploader
@@ -133,20 +157,6 @@ steps:
        env:
          DOCKER_BUILDKIT: "1"

-      - label: "Build wheel - x86_64 - CUDA 13.0"
-        depends_on: ~
-        id: build-wheel-x86-cuda-13-0
-        agents:
-          queue: cpu_queue_release
-        commands:
-          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=13.0.2 --build-arg torch_cuda_arch_list=\"${CUDA_ARCH_X86}\" --build-arg BUILD_OS=manylinux --build-arg BUILD_BASE_IMAGE=pytorch/manylinux2_28-builder:cuda13.0 --tag vllm-ci:build-image --target build --progress plain -f docker/Dockerfile ."
-          - "mkdir artifacts"
-          - "docker run --rm -v $(pwd)/artifacts:/artifacts_host vllm-ci:build-image bash -c 'cp -r dist /artifacts_host && chmod -R a+rw /artifacts_host'"
-          - "bash .buildkite/scripts/upload-nightly-wheels.sh"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "s3://vllm-wheels/$$BUILDKITE_COMMIT/$(cd artifacts/dist && echo *.whl)" release-wheels'
-        env:
-          DOCKER_BUILDKIT: "1"
-
      - label: "Build wheel - x86_64 - CPU"
        depends_on: ~
        id: build-wheel-x86-cpu
@@ -162,12 +172,26 @@ steps:
          DOCKER_BUILDKIT: "1"

  - label: "Generate and upload wheel indices"
+    key: generate-wheel-indices
    depends_on: "build-wheels"
    allow_dependency_failure: true
+    if: build.env("NIGHTLY") != "1"
+    agents:
+      queue: cpu_queue_release
+    commands:
+      - "UPDATE_VERSION_INDEX=0 bash .buildkite/scripts/generate-and-upload-nightly-index.sh"
+
+  - label: "Regenerate indices with additional wheels"
+    key: generate-additional-wheel-indices
+    depends_on:
+      - build-wheels
+      - build-additional-wheels
+      - generate-wheel-indices
+    allow_dependency_failure: true
    agents:
      queue: cpu_queue_release
    commands:
-      - "bash .buildkite/scripts/generate-and-upload-nightly-index.sh"
+      - 'UPDATE_NIGHTLY_INDEX="$${NIGHTLY:-0}" bash .buildkite/scripts/generate-and-upload-nightly-index.sh'

  - block: "Unblock to build release Docker images"
    depends_on: ~
@@ -566,366 +590,370 @@ steps:
  #
  # =============================================================================

-  # ROCm Job 1: Build ROCm Base Wheels (with S3 caching)
-  - label: ":rocm: Build ROCm Base Image & Wheels"
-    id: build-rocm-base-wheels
+  - group: "Build ROCm Wheel / Image "
+    key: "build-rocm-wheel-image"
    depends_on: ~
-    agents:
-      queue: cpu_queue_release
-    commands:
-      - |
-        set -euo pipefail
+    steps:
+      # ROCm Job 1: Build ROCm Base Wheels (with S3 caching)
+      - label: ":rocm: Build ROCm Base Image & Wheels"
+        id: build-rocm-base-wheels
+        depends_on: ~
+        agents:
+          queue: cpu_queue_release
+        commands:
+          - |
+            set -euo pipefail

-        # Generate cache key
-        CACHE_KEY=$$(.buildkite/scripts/cache-rocm-base-wheels.sh key)
-        ECR_CACHE_TAG="public.ecr.aws/q9t5s3a7/vllm-release-repo:$${CACHE_KEY}-rocm-base"
+            # Generate cache key
+            CACHE_KEY=$$(.buildkite/scripts/cache-rocm-base-wheels.sh key)
+            ECR_CACHE_TAG="public.ecr.aws/q9t5s3a7/vllm-release-repo:$${CACHE_KEY}-rocm-base"

-        echo "========================================"
-        echo "ROCm Base Build Configuration"
-        echo "========================================"
-        echo "  CACHE_KEY: $${CACHE_KEY}"
-        echo "  ECR_CACHE_TAG: $${ECR_CACHE_TAG}"
-        echo "========================================"
-        
-        # Login to ECR
-        aws ecr-public get-login-password --region us-east-1 | \
-          docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7
-        
-        IMAGE_EXISTS=false
-        WHEELS_EXIST=false
-        
-        # Check ECR for Docker image
+            echo "========================================"
+            echo "ROCm Base Build Configuration"
+            echo "========================================"
+            echo "  CACHE_KEY: $${CACHE_KEY}"
+            echo "  ECR_CACHE_TAG: $${ECR_CACHE_TAG}"
+            echo "========================================"
+            
+            # Login to ECR
+            aws ecr-public get-login-password --region us-east-1 | \
+              docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7
+            
+            IMAGE_EXISTS=false
+            WHEELS_EXIST=false
+            
+            # Check ECR for Docker image

-        if docker manifest inspect "$${ECR_CACHE_TAG}" > /dev/null 2>&1; then
-          IMAGE_EXISTS=true
-          echo "ECR image cache HIT"
-        fi
-        
-        # Check S3 for wheels
-        WHEEL_CACHE_STATUS=$(.buildkite/scripts/cache-rocm-base-wheels.sh check)
-        if [ "$${WHEEL_CACHE_STATUS}" = "hit" ]; then
-          WHEELS_EXIST=true
-          echo "S3 wheels cache HIT"
-        fi
+            if docker manifest inspect "$${ECR_CACHE_TAG}" > /dev/null 2>&1; then
+              IMAGE_EXISTS=true
+              echo "ECR image cache HIT"
+            fi
+            
+            # Check S3 for wheels
+            WHEEL_CACHE_STATUS=$(.buildkite/scripts/cache-rocm-base-wheels.sh check)
+            if [ "$${WHEEL_CACHE_STATUS}" = "hit" ]; then
+              WHEELS_EXIST=true
+              echo "S3 wheels cache HIT"
+            fi

-        
-        # Scenario 1: Both cached (best case)
-        if [ "$${IMAGE_EXISTS}" = "true" ] && [ "$${WHEELS_EXIST}" = "true" ]; then
-          echo ""
-          echo "FULL CACHE HIT - Reusing both image and wheels"
-          echo ""
+            
+            # Scenario 1: Both cached (best case)
+            if [ "$${IMAGE_EXISTS}" = "true" ] && [ "$${WHEELS_EXIST}" = "true" ]; then
+              echo ""
+              echo "FULL CACHE HIT - Reusing both image and wheels"
+              echo ""

-          # Download wheels
-          .buildkite/scripts/cache-rocm-base-wheels.sh download
-          
-          # Save ECR tag for downstream jobs
-          buildkite-agent meta-data set "rocm-base-image-tag" "$${ECR_CACHE_TAG}"
-          
-        # Scenario 2: Full rebuild needed
-        else
-          echo ""
-          echo " CACHE MISS - Building from scratch..."
-          echo ""
-          
-          # Build full base image and push to ECR
-          DOCKER_BUILDKIT=1 docker buildx build \
-            --file docker/Dockerfile.rocm_base \
-            --tag "$${ECR_CACHE_TAG}" \
-            --build-arg USE_SCCACHE=1 \
-            --build-arg SCCACHE_BUCKET_NAME=vllm-build-sccache \
-            --build-arg SCCACHE_REGION_NAME=us-west-2 \
-            --build-arg SCCACHE_S3_NO_CREDENTIALS=0 \
-            --push \
-            .
-          
-          # Build wheel extraction stage
-          DOCKER_BUILDKIT=1 docker buildx build \
-            --file docker/Dockerfile.rocm_base \
-            --tag rocm-base-debs:$${BUILDKITE_BUILD_NUMBER} \
-            --target debs_wheel_release \
-            --build-arg USE_SCCACHE=1 \
-            --build-arg SCCACHE_BUCKET_NAME=vllm-build-sccache \
-            --build-arg SCCACHE_REGION_NAME=us-west-2 \
-            --build-arg SCCACHE_S3_NO_CREDENTIALS=0 \
-            --load \
-            .
-          
-          # Extract and upload wheels
-          mkdir -p artifacts/rocm-base-wheels
-          cid=$(docker create rocm-base-debs:$${BUILDKITE_BUILD_NUMBER})
-          docker cp $${cid}:/app/debs/. artifacts/rocm-base-wheels/
-          docker rm $${cid}
-          
-          .buildkite/scripts/cache-rocm-base-wheels.sh upload
+              # Download wheels
+              .buildkite/scripts/cache-rocm-base-wheels.sh download
+              
+              # Save ECR tag for downstream jobs
+              buildkite-agent meta-data set "rocm-base-image-tag" "$${ECR_CACHE_TAG}"
+              
+            # Scenario 2: Full rebuild needed
+            else
+              echo ""
+              echo " CACHE MISS - Building from scratch..."
+              echo ""
+              
+              # Build full base image and push to ECR
+              DOCKER_BUILDKIT=1 docker buildx build \
+                --file docker/Dockerfile.rocm_base \
+                --tag "$${ECR_CACHE_TAG}" \
+                --build-arg USE_SCCACHE=1 \
+                --build-arg SCCACHE_BUCKET_NAME=vllm-build-sccache \
+                --build-arg SCCACHE_REGION_NAME=us-west-2 \
+                --build-arg SCCACHE_S3_NO_CREDENTIALS=0 \
+                --push \
+                .
+              
+              # Build wheel extraction stage
+              DOCKER_BUILDKIT=1 docker buildx build \
+                --file docker/Dockerfile.rocm_base \
+                --tag rocm-base-debs:$${BUILDKITE_BUILD_NUMBER} \
+                --target debs_wheel_release \
+                --build-arg USE_SCCACHE=1 \
+                --build-arg SCCACHE_BUCKET_NAME=vllm-build-sccache \
+                --build-arg SCCACHE_REGION_NAME=us-west-2 \
+                --build-arg SCCACHE_S3_NO_CREDENTIALS=0 \
+                --load \
+                .
+              
+              # Extract and upload wheels
+              mkdir -p artifacts/rocm-base-wheels
+              cid=$(docker create rocm-base-debs:$${BUILDKITE_BUILD_NUMBER})
+              docker cp $${cid}:/app/debs/. artifacts/rocm-base-wheels/
+              docker rm $${cid}
+              
+              .buildkite/scripts/cache-rocm-base-wheels.sh upload

-          # Cache base docker image to ECR
-          docker push "$${ECR_CACHE_TAG}"
-          
-          buildkite-agent meta-data set "rocm-base-image-tag" "$${ECR_CACHE_TAG}"
-          
-          echo ""
-          echo " Build complete - Image and wheels cached"
-        fi
+              # Cache base docker image to ECR
+              docker push "$${ECR_CACHE_TAG}"
+              
+              buildkite-agent meta-data set "rocm-base-image-tag" "$${ECR_CACHE_TAG}"
+              
+              echo ""
+              echo " Build complete - Image and wheels cached"
+            fi

-    artifact_paths:
-      - "artifacts/rocm-base-wheels/*.whl"
-    env:
-      DOCKER_BUILDKIT: "1"
-      S3_BUCKET: "vllm-wheels"
+        artifact_paths:
+          - "artifacts/rocm-base-wheels/*.whl"
+        env:
+          DOCKER_BUILDKIT: "1"
+          S3_BUCKET: "vllm-wheels"

-  # ROCm Job 2: Build vLLM ROCm Wheel
-  - label: ":python: Build vLLM ROCm Wheel - x86_64"
-    id: build-rocm-vllm-wheel
-    depends_on:
-      - step: build-rocm-base-wheels
-        allow_failure: false
-    agents:
-      queue: cpu_queue_release
-    timeout_in_minutes: 180
-    commands:
-      # Download artifacts and prepare Docker image
-      - |
-        set -euo pipefail
+      # ROCm Job 2: Build vLLM ROCm Wheel
+      - label: ":python: Build vLLM ROCm Wheel - x86_64"
+        id: build-rocm-vllm-wheel
+        depends_on:
+          - step: build-rocm-base-wheels
+            allow_failure: false
+        agents:
+          queue: cpu_queue_release
+        timeout_in_minutes: 180
+        commands:
+          # Download artifacts and prepare Docker image
+          - |
+            set -euo pipefail

-        # Ensure git tags are up-to-date (Buildkite's default fetch doesn't update tags)
-        # This fixes version detection when tags are moved/force-pushed
-        echo "Fetching latest tags from origin..."
-        git fetch --tags --force origin
-        
-        # Log tag information for debugging version detection
-        echo "========================================"
-        echo "Git Tag Verification"
-        echo "========================================"
-        echo "Current HEAD: $(git rev-parse HEAD)"
-        echo "git describe --tags: $(git describe --tags 2>/dev/null || echo 'No tags found')"
-        echo ""
-        echo "Recent tags (pointing to commits near HEAD):"
-        git tag -l --sort=-creatordate | head -5
-        echo "setuptools_scm version detection:"
-        pip install -q setuptools_scm 2>/dev/null || true
-        python3 -c "import setuptools_scm; print('  Detected version:', setuptools_scm.get_version())" 2>/dev/null || echo "  (setuptools_scm not available in this environment)"
-        echo "========================================"
+            # Ensure git tags are up-to-date (Buildkite's default fetch doesn't update tags)
+            # This fixes version detection when tags are moved/force-pushed
+            echo "Fetching latest tags from origin..."
+            git fetch --tags --force origin
+            
+            # Log tag information for debugging version detection
+            echo "========================================"
+            echo "Git Tag Verification"
+            echo "========================================"
+            echo "Current HEAD: $(git rev-parse HEAD)"
+            echo "git describe --tags: $(git describe --tags 2>/dev/null || echo 'No tags found')"
+            echo ""
+            echo "Recent tags (pointing to commits near HEAD):"
+            git tag -l --sort=-creatordate | head -5
+            echo "setuptools_scm version detection:"
+            pip install -q setuptools_scm 2>/dev/null || true
+            python3 -c "import setuptools_scm; print('  Detected version:', setuptools_scm.get_version())" 2>/dev/null || echo "  (setuptools_scm not available in this environment)"
+            echo "========================================"

-        # Download wheel artifacts from current build
-        echo "Downloading wheel artifacts from current build"
-        buildkite-agent artifact download "artifacts/rocm-base-wheels/*.whl" .
+            # Download wheel artifacts from current build
+            echo "Downloading wheel artifacts from current build"
+            buildkite-agent artifact download "artifacts/rocm-base-wheels/*.whl" .

-        # Get ECR image tag from metadata (set by build-rocm-base-wheels)
-        ECR_IMAGE_TAG="$$(buildkite-agent meta-data get rocm-base-image-tag 2>/dev/null || echo '')"
-        if [ -z "$${ECR_IMAGE_TAG}" ]; then
-          echo "ERROR: rocm-base-image-tag metadata not found"
-          echo "This should have been set by the build-rocm-base-wheels job"
-          exit 1
-        fi
-        
-        echo "Pulling base Docker image from ECR: $${ECR_IMAGE_TAG}"
-        
-        # Login to ECR
-        aws ecr-public get-login-password --region us-east-1 | \
-          docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7
-        
-        # Pull base Docker image from ECR
-        docker pull "$${ECR_IMAGE_TAG}"
-        
-        echo "Loaded base image: $${ECR_IMAGE_TAG}"
-        
-        # Prepare base wheels for Docker build context
-        mkdir -p docker/context/base-wheels
-        touch docker/context/base-wheels/.keep
-        cp artifacts/rocm-base-wheels/*.whl docker/context/base-wheels/
-        echo "Base wheels for vLLM build:"
-        ls -lh docker/context/base-wheels/
+            # Get ECR image tag from metadata (set by build-rocm-base-wheels)
+            ECR_IMAGE_TAG="$$(buildkite-agent meta-data get rocm-base-image-tag 2>/dev/null || echo '')"
+            if [ -z "$${ECR_IMAGE_TAG}" ]; then
+              echo "ERROR: rocm-base-image-tag metadata not found"
+              echo "This should have been set by the build-rocm-base-wheels job"
+              exit 1
+            fi
+            
+            echo "Pulling base Docker image from ECR: $${ECR_IMAGE_TAG}"
+            
+            # Login to ECR
+            aws ecr-public get-login-password --region us-east-1 | \
+              docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7
+            
+            # Pull base Docker image from ECR
+            docker pull "$${ECR_IMAGE_TAG}"
+            
+            echo "Loaded base image: $${ECR_IMAGE_TAG}"
+            
+            # Prepare base wheels for Docker build context
+            mkdir -p docker/context/base-wheels
+            touch docker/context/base-wheels/.keep
+            cp artifacts/rocm-base-wheels/*.whl docker/context/base-wheels/
+            echo "Base wheels for vLLM build:"
+            ls -lh docker/context/base-wheels/

-        echo "========================================"
-        echo "Building vLLM wheel with:"
-        echo "  BUILDKITE_COMMIT: $${BUILDKITE_COMMIT}"
-        echo "  BUILDKITE_BRANCH: $${BUILDKITE_BRANCH}"
-        echo "  BASE_IMAGE: $${ECR_IMAGE_TAG}"
-        echo "========================================"
+            echo "========================================"
+            echo "Building vLLM wheel with:"
+            echo "  BUILDKITE_COMMIT: $${BUILDKITE_COMMIT}"
+            echo "  BUILDKITE_BRANCH: $${BUILDKITE_BRANCH}"
+            echo "  BASE_IMAGE: $${ECR_IMAGE_TAG}"
+            echo "========================================"

-        # Build vLLM wheel using local checkout (REMOTE_VLLM=0)
-        DOCKER_BUILDKIT=1 docker build \
-          --file docker/Dockerfile.rocm \
-          --target export_vllm_wheel_release \
-          --output type=local,dest=rocm-dist \
-          --build-arg BASE_IMAGE="$${ECR_IMAGE_TAG}" \
-          --build-arg REMOTE_VLLM=0 \
-          --build-arg GIT_REPO_CHECK=1 \
-          --build-arg USE_SCCACHE=1 \
-          --build-arg SCCACHE_BUCKET_NAME=vllm-build-sccache \
-          --build-arg SCCACHE_REGION_NAME=us-west-2 \
-          --build-arg SCCACHE_S3_NO_CREDENTIALS=0 \
-          .
-        echo "Built vLLM wheel:"
-        ls -lh rocm-dist/*.whl
-        # Copy wheel to artifacts directory
-        mkdir -p artifacts/rocm-vllm-wheel
-        cp rocm-dist/*.whl artifacts/rocm-vllm-wheel/
-        echo "Final vLLM wheel:"
-        ls -lh artifacts/rocm-vllm-wheel/
-    artifact_paths:
-      - "artifacts/rocm-vllm-wheel/*.whl"
-    env:
-      DOCKER_BUILDKIT: "1"
-      S3_BUCKET: "vllm-wheels"
+            # Build vLLM wheel using local checkout (REMOTE_VLLM=0)
+            DOCKER_BUILDKIT=1 docker build \
+              --file docker/Dockerfile.rocm \
+              --target export_vllm_wheel_release \
+              --output type=local,dest=rocm-dist \
+              --build-arg BASE_IMAGE="$${ECR_IMAGE_TAG}" \
+              --build-arg REMOTE_VLLM=0 \
+              --build-arg GIT_REPO_CHECK=1 \
+              --build-arg USE_SCCACHE=1 \
+              --build-arg SCCACHE_BUCKET_NAME=vllm-build-sccache \
+              --build-arg SCCACHE_REGION_NAME=us-west-2 \
+              --build-arg SCCACHE_S3_NO_CREDENTIALS=0 \
+              .
+            echo "Built vLLM wheel:"
+            ls -lh rocm-dist/*.whl
+            # Copy wheel to artifacts directory
+            mkdir -p artifacts/rocm-vllm-wheel
+            cp rocm-dist/*.whl artifacts/rocm-vllm-wheel/
+            echo "Final vLLM wheel:"
+            ls -lh artifacts/rocm-vllm-wheel/
+        artifact_paths:
+          - "artifacts/rocm-vllm-wheel/*.whl"
+        env:
+          DOCKER_BUILDKIT: "1"
+          S3_BUCKET: "vllm-wheels"

-  # ROCm Job 3: Upload Wheels to S3
-  - label: ":s3: Upload ROCm Wheels to S3"
-    id: upload-rocm-wheels
-    depends_on:
-      - step: build-rocm-vllm-wheel
-        allow_failure: false
-    agents:
-      queue: cpu_queue_release
-    timeout_in_minutes: 60
-    commands:
-      # Download all wheel artifacts and run upload
-      - |
-        set -euo pipefail
+      # ROCm Job 3: Upload Wheels to S3
+      - label: ":s3: Upload ROCm Wheels to S3"
+        id: upload-rocm-wheels
+        depends_on:
+          - step: build-rocm-vllm-wheel
+            allow_failure: false
+        agents:
+          queue: cpu_queue_release
+        timeout_in_minutes: 60
+        commands:
+          # Download all wheel artifacts and run upload
+          - |
+            set -euo pipefail

-        # Download artifacts from current build
-        echo "Downloading artifacts from current build"
-        buildkite-agent artifact download "artifacts/rocm-base-wheels/*.whl" .
-        buildkite-agent artifact download "artifacts/rocm-vllm-wheel/*.whl" .
+            # Download artifacts from current build
+            echo "Downloading artifacts from current build"
+            # buildkite-agent artifact download "artifacts/rocm-base-wheels/*.whl" .
+            # buildkite-agent artifact download "artifacts/rocm-vllm-wheel/*.whl" .

-        # Run upload script
-        bash .buildkite/scripts/upload-rocm-wheels.sh
-    env:
-      DOCKER_BUILDKIT: "1"
-      S3_BUCKET: "vllm-wheels"
+            # # Run upload script
+            bash .buildkite/scripts/upload-rocm-wheels.sh
+        env:
+          DOCKER_BUILDKIT: "1"
+          S3_BUCKET: "vllm-wheels"

-  # ROCm Job 4: Annotate ROCm Wheel Release
-  - label: ":memo: Annotate ROCm wheel release"
-    id: annotate-rocm-release
-    depends_on:
-      - upload-rocm-wheels
-    agents:
-      queue: cpu_queue_release
-    commands:
-      - "bash .buildkite/scripts/annotate-rocm-release.sh"
-    env:
-      S3_BUCKET: "vllm-wheels"
+      # ROCm Job 4: Annotate ROCm Wheel Release
+      - label: ":memo: Annotate ROCm wheel release"
+        id: annotate-rocm-release
+        depends_on:
+          - upload-rocm-wheels
+        agents:
+          queue: cpu_queue_release
+        commands:
+          - "bash .buildkite/scripts/annotate-rocm-release.sh"
+        env:
+          S3_BUCKET: "vllm-wheels"

-  # ROCm Job 5: Generate Root Index for ROCm Wheels (for release only)
-  # This is the job to create https://wheels.vllm.ai/rocm/ index allowing
-  # users to install with `uv pip install vllm --extra-index-url https://wheels.vllm.ai/rocm/`
-  - block: "Generate Root Index for ROCm Wheels for Release"
-    key: block-generate-root-index-rocm-wheels
-    depends_on: upload-rocm-wheels
+      # ROCm Job 5: Generate Root Index for ROCm Wheels (for release only)
+      # This is the job to create https://wheels.vllm.ai/rocm/ index allowing
+      # users to install with `uv pip install vllm --extra-index-url https://wheels.vllm.ai/rocm/`
+      - block: "Generate Root Index for ROCm Wheels for Release"
+        key: block-generate-root-index-rocm-wheels
+        depends_on: upload-rocm-wheels

-  - label: ":package: Generate Root Index for ROCm Wheels for Release"
-    depends_on: block-generate-root-index-rocm-wheels
-    id: generate-root-index-rocm-wheels
-    agents:
-      queue: cpu_queue_release
-    commands:
-      - "bash tools/vllm-rocm/generate-rocm-wheels-root-index.sh"
-    env:
-      S3_BUCKET: "vllm-wheels"
-      VARIANT: "rocm723"
+      - label: ":package: Generate Root Index for ROCm Wheels for Release"
+        depends_on: block-generate-root-index-rocm-wheels
+        id: generate-root-index-rocm-wheels
+        agents:
+          queue: cpu_queue_release
+        commands:
+          - "bash tools/vllm-rocm/generate-rocm-wheels-root-index.sh"
+        env:
+          S3_BUCKET: "vllm-wheels"
+          VARIANT: "rocm723"

-  # ROCm Job 6: Build ROCm Release Docker Image
-  - label: ":docker: Build release image - x86_64 - ROCm"
-    id: build-rocm-release-image
-    depends_on:
-      - step: block-build-release-images
-        allow_failure: true
-      - step: build-rocm-base-wheels
-        allow_failure: false
-    agents:
-      queue: cpu_queue_release
-    timeout_in_minutes: 60
-    commands:
-      - |
-        set -euo pipefail
-        
-        # Login to ECR
-        aws ecr-public get-login-password --region us-east-1 | \
-          docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7
-        
-        # Get ECR image tag from metadata (set by build-rocm-base-wheels)
-        ECR_IMAGE_TAG="$$(buildkite-agent meta-data get rocm-base-image-tag 2>/dev/null || echo '')"
-        if [ -z "$${ECR_IMAGE_TAG}" ]; then
-          echo "ERROR: rocm-base-image-tag metadata not found"
-          echo "This should have been set by the build-rocm-base-wheels job"
-          exit 1
-        fi
-        
-        echo "Pulling base Docker image from ECR: $${ECR_IMAGE_TAG}"
-        
-        # Pull base Docker image from ECR
-        docker pull "$${ECR_IMAGE_TAG}"
-        
-        echo "Loaded base image: $${ECR_IMAGE_TAG}"
-        
-        # Pass the base image ECR tag to downstream steps (nightly publish)
-        buildkite-agent meta-data set "rocm-base-ecr-tag" "$${ECR_IMAGE_TAG}"
-        
-        echo "========================================"
-        echo "Building vLLM ROCm release image with:"
-        echo "  BASE_IMAGE: $${ECR_IMAGE_TAG}"
-        echo "  BUILDKITE_COMMIT: $${BUILDKITE_COMMIT}"
-        echo "========================================"
-        
-        # Build vLLM ROCm release image using cached base
-        DOCKER_BUILDKIT=1 docker build \
-          --build-arg max_jobs=16 \
-          --build-arg BASE_IMAGE="$${ECR_IMAGE_TAG}" \
-          --build-arg USE_SCCACHE=1 \
-          --build-arg SCCACHE_BUCKET_NAME=vllm-build-sccache \
-          --build-arg SCCACHE_REGION_NAME=us-west-2 \
-          --build-arg SCCACHE_S3_NO_CREDENTIALS=0 \
-          --tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$${BUILDKITE_COMMIT}-rocm \
-          --target vllm-openai \
-          --progress plain \
-          -f docker/Dockerfile.rocm .
-        
-        # Push to ECR
-        docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$${BUILDKITE_COMMIT}-rocm
+      # ROCm Job 6: Build ROCm Release Docker Image
+      - label: ":docker: Build release image - x86_64 - ROCm"
+        id: build-rocm-release-image
+        depends_on:
+          - step: block-build-release-images
+            allow_failure: true
+          - step: build-rocm-base-wheels
+            allow_failure: false
+        agents:
+          queue: cpu_queue_release
+        timeout_in_minutes: 60
+        commands:
+          - |
+            set -euo pipefail
+            
+            # Login to ECR
+            aws ecr-public get-login-password --region us-east-1 | \
+              docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7
+            
+            # Get ECR image tag from metadata (set by build-rocm-base-wheels)
+            ECR_IMAGE_TAG="$$(buildkite-agent meta-data get rocm-base-image-tag 2>/dev/null || echo '')"
+            if [ -z "$${ECR_IMAGE_TAG}" ]; then
+              echo "ERROR: rocm-base-image-tag metadata not found"
+              echo "This should have been set by the build-rocm-base-wheels job"
+              exit 1
+            fi
+            
+            echo "Pulling base Docker image from ECR: $${ECR_IMAGE_TAG}"
+            
+            # Pull base Docker image from ECR
+            docker pull "$${ECR_IMAGE_TAG}"
+            
+            echo "Loaded base image: $${ECR_IMAGE_TAG}"
+            
+            # Pass the base image ECR tag to downstream steps (nightly publish)
+            buildkite-agent meta-data set "rocm-base-ecr-tag" "$${ECR_IMAGE_TAG}"
+            
+            echo "========================================"
+            echo "Building vLLM ROCm release image with:"
+            echo "  BASE_IMAGE: $${ECR_IMAGE_TAG}"
+            echo "  BUILDKITE_COMMIT: $${BUILDKITE_COMMIT}"
+            echo "========================================"
+            
+            # Build vLLM ROCm release image using cached base
+            DOCKER_BUILDKIT=1 docker build \
+              --build-arg max_jobs=16 \
+              --build-arg BASE_IMAGE="$${ECR_IMAGE_TAG}" \
+              --build-arg USE_SCCACHE=1 \
+              --build-arg SCCACHE_BUCKET_NAME=vllm-build-sccache \
+              --build-arg SCCACHE_REGION_NAME=us-west-2 \
+              --build-arg SCCACHE_S3_NO_CREDENTIALS=0 \
+              --tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$${BUILDKITE_COMMIT}-rocm \
+              --target vllm-openai \
+              --progress plain \
+              -f docker/Dockerfile.rocm .
+            
+            # Push to ECR
+            docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$${BUILDKITE_COMMIT}-rocm

-        echo ""
-        echo " Successfully built and pushed ROCm release image"
-        echo "   Image: public.ecr.aws/q9t5s3a7/vllm-release-repo:$${BUILDKITE_COMMIT}-rocm"
-        echo ""
-    env:
-      DOCKER_BUILDKIT: "1"
-      S3_BUCKET: "vllm-wheels"
+            echo ""
+            echo " Successfully built and pushed ROCm release image"
+            echo "   Image: public.ecr.aws/q9t5s3a7/vllm-release-repo:$${BUILDKITE_COMMIT}-rocm"
+            echo ""
+        env:
+          DOCKER_BUILDKIT: "1"
+          S3_BUCKET: "vllm-wheels"

-  - label: "Publish nightly XPU image to DockerHub"
-    depends_on:
-      - create-manifest-xpu
-    if: build.env("NIGHTLY") == "1"
-    agents:
-      queue: small_cpu_queue_release
-    commands:
-      - "bash .buildkite/scripts/xpu/push-nightly-builds-xpu.sh"
-      - "bash .buildkite/scripts/cleanup-nightly-builds.sh nightly- vllm/vllm-openai-xpu"
-    plugins:
-      - docker-login#v3.0.0:
-          username: vllmbot
-          password-env: DOCKERHUB_TOKEN
-    env:
-      DOCKER_BUILDKIT: "1"
-      DOCKERHUB_USERNAME: "vllmbot"
+      - label: "Publish nightly XPU image to DockerHub"
+        depends_on:
+          - create-manifest-xpu
+        if: build.env("NIGHTLY") == "1"
+        agents:
+          queue: small_cpu_queue_release
+        commands:
+          - "bash .buildkite/scripts/xpu/push-nightly-builds-xpu.sh"
+          - "bash .buildkite/scripts/cleanup-nightly-builds.sh nightly- vllm/vllm-openai-xpu"
+        plugins:
+          - docker-login#v3.0.0:
+              username: vllmbot
+              password-env: DOCKERHUB_TOKEN
+        env:
+          DOCKER_BUILDKIT: "1"
+          DOCKERHUB_USERNAME: "vllmbot"

-  - label: "Publish nightly ROCm image to DockerHub"
-    depends_on:
-      - build-rocm-release-image
-    if: build.env("NIGHTLY") == "1"
-    agents:
-      queue: small_cpu_queue_release
-    commands:
-      - "bash .buildkite/scripts/push-nightly-builds-rocm.sh"
-      # Clean up old nightly builds (keep only last 14)
-      - "bash .buildkite/scripts/cleanup-nightly-builds.sh nightly- vllm/vllm-openai-rocm"
-      - "bash .buildkite/scripts/cleanup-nightly-builds.sh base-nightly- vllm/vllm-openai-rocm"
-    plugins:
-      - docker-login#v3.0.0:
-          username: vllmbot
-          password-env: DOCKERHUB_TOKEN
-    env:
-      DOCKER_BUILDKIT: "1"
-      DOCKERHUB_USERNAME: "vllmbot"
+      - label: "Publish nightly ROCm image to DockerHub"
+        depends_on:
+          - build-rocm-release-image
+        if: build.env("NIGHTLY") == "1"
+        agents:
+          queue: small_cpu_queue_release
+        commands:
+          - "bash .buildkite/scripts/push-nightly-builds-rocm.sh"
+          # Clean up old nightly builds (keep only last 14)
+          - "bash .buildkite/scripts/cleanup-nightly-builds.sh nightly- vllm/vllm-openai-rocm"
+          - "bash .buildkite/scripts/cleanup-nightly-builds.sh base-nightly- vllm/vllm-openai-rocm"
+        plugins:
+          - docker-login#v3.0.0:
+              username: vllmbot
+              password-env: DOCKERHUB_TOKEN
+        env:
+          DOCKER_BUILDKIT: "1"
+          DOCKERHUB_USERNAME: "vllmbot"

  # =============================================================================
  # Publish to DockerHub and PyPI (at the end so all builds complete first)
@@ -974,6 +1002,8 @@ steps:
        depends_on:
          - input-release-version
          - build-wheels
+          - build-additional-wheels
+          - generate-additional-wheel-indices

      - label: "Upload release wheels to PyPI"
        depends_on:
@@ -45,8 +45,10 @@ $PYTHON .buildkite/scripts/generate-nightly-index.py --version "$SUBPATH" --curr
 echo "Uploading indices to $S3_COMMIT_PREFIX"
 aws s3 cp --recursive "$INDICES_OUTPUT_DIR/" "$S3_COMMIT_PREFIX"

-# copy to /nightly/ only if it is on the main branch and not a PR
-if [[ "$BUILDKITE_BRANCH" == "main" && "$BUILDKITE_PULL_REQUEST" == "false" ]]; then
+# copy to /nightly/ only when enabled for a main branch build that is not a PR
+if [[ "${UPDATE_NIGHTLY_INDEX:-1}" == "1" && \
+      "$BUILDKITE_BRANCH" == "main" && \
+      "$BUILDKITE_PULL_REQUEST" == "false" ]]; then
    echo "Uploading indices to overwrite /nightly/"
    aws s3 cp --recursive "$INDICES_OUTPUT_DIR/" "s3://$BUCKET/nightly/"
 fi
@@ -67,7 +69,7 @@ pure_version="${version%%+*}"
 echo "Pure version (without variant): $pure_version"

 # re-generate and copy to /<pure_version>/ only if it does not have "dev" in the version
-if [[ "$version" != *"dev"* ]]; then
+if [[ "${UPDATE_VERSION_INDEX:-1}" == "1" && "$version" != *"dev"* ]]; then
    echo "Re-generating indices for /$pure_version/"
    rm -rf "${INDICES_OUTPUT_DIR:?}"
    mkdir -p "$INDICES_OUTPUT_DIR"
@@ -40,7 +40,9 @@ function cpu_tests() {
    pytest -x -v -s tests/kernels/moe/test_cpu_fused_moe.py
    pytest -x -v -s tests/kernels/mamba/cpu/test_cpu_gdn_ops.py
    pytest -x -v -s tests/kernels/moe/test_cpu_int4_moe.py
-    pytest -x -v -s tests/kernels/mamba/test_cpu_short_conv.py"
+    pytest -x -v -s tests/kernels/mamba/test_cpu_short_conv.py
+    pytest -x -v -s tests/kernels/mamba/test_causal_conv1d.py
+    pytest -x -v -s tests/kernels/mamba/test_mamba_ssm.py"

  # skip tests requiring model downloads if HF_TOKEN is not set
  # due to rate-limits
@@ -97,3 +99,4 @@ function cpu_tests() {
 # All of CPU tests are expected to be finished less than 40 mins.
 export -f cpu_tests
 timeout 2h bash -c cpu_tests
+
@@ -113,8 +113,8 @@ $PYTHON .buildkite/scripts/generate-nightly-index.py \
 echo "Uploading indices to $S3_COMMIT_PREFIX"
 aws s3 cp --recursive "$INDICES_OUTPUT_DIR/" "$S3_COMMIT_PREFIX"

-# Update rocm/nightly/ if on main branch and not a PR
-if [[ "$BUILDKITE_BRANCH" == "main" && "$BUILDKITE_PULL_REQUEST" == "false" ]] || [[ "$NIGHTLY" == "1" ]]; then
+# Only scheduled nightly builds should update the moving nightly index.
+if [[ "${NIGHTLY:-0}" == "1" ]]; then
    echo "Updating rocm/nightly/ index..."
    aws s3 cp --recursive "$INDICES_OUTPUT_DIR/" "s3://$BUCKET/rocm/nightly/"
 fi
@@ -147,7 +147,7 @@ echo ""
 echo "Install command (by commit):"
 echo "  pip install vllm --extra-index-url https://${BUCKET}.s3.amazonaws.com/$ROCM_SUBPATH/"
 echo ""
-if [[ "$BUILDKITE_BRANCH" == "main" ]] || [[ "$NIGHTLY" == "1" ]]; then
+if [[ "${NIGHTLY:-0}" == "1" ]]; then
    echo "Install command (nightly):"
    echo "  pip install vllm --extra-index-url https://${BUCKET}.s3.amazonaws.com/rocm/nightly/"
 fi
@@ -18,6 +18,7 @@ steps:
    - pytest -v -s cuda/test_platform_no_cuda_init.py

 - label: Cudagraph
+  device: h200_35gb
  key: cudagraph
  timeout_in_minutes: 30
  source_file_dependencies:
@@ -28,4 +29,4 @@ steps:
  commands:
    - pytest -v -s v1/cudagraph/test_cudagraph_dispatch.py
    - pytest -v -s v1/cudagraph/test_cudagraph_mode.py
-    - pytest -v -s v1/cudagraph/test_breakable_cudagraph.py
+    - pytest -v -s v1/cudagraph/test_breakable_cudagraph.py
@@ -3,6 +3,7 @@ depends_on:
  - image-build
 steps:
 - label: Entrypoints Unit Tests
+  device: h200_35gb
  key: entrypoints-unit-tests
  timeout_in_minutes: 25
  working_dir: "/vllm-workspace/tests"
@@ -15,6 +16,7 @@ steps:
  - pytest -v -s entrypoints/weight_transfer

 - label: Entrypoints Integration (LLM)
+  device: h200_35gb
  key: entrypoints-integration-llm
  timeout_in_minutes: 60
  working_dir: "/vllm-workspace/tests"
@@ -55,6 +57,7 @@ steps:
      - image-build-amd

 - label: Entrypoints Integration (API Server OpenAI - Part 1)
+  device: h200_35gb
  key: entrypoints-integration-api-server-openai-part-1
  timeout_in_minutes: 45
  working_dir: "/vllm-workspace/tests"
@@ -73,6 +76,7 @@ steps:
      - image-build-amd

 - label: Entrypoints Integration (API Server OpenAI - Part 2)
+  device: h200_35gb
  key: entrypoints-integration-api-server-openai-part-2
  timeout_in_minutes: 45
  working_dir: "/vllm-workspace/tests"
@@ -92,6 +96,7 @@ steps:
      - image-build-amd

 - label: Entrypoints Integration (API Server Generate)
+  device: h200_35gb
  key: entrypoints-integration-api-server-generate
  timeout_in_minutes: 50
  working_dir: "/vllm-workspace/tests"
@@ -114,6 +119,7 @@ steps:
      - image-build-amd

 - label: Entrypoints Integration (Responses API)
+  device: h200_35gb
  key: entrypoints-integration-responses-api
  timeout_in_minutes: 50
  working_dir: "/vllm-workspace/tests"
@@ -148,6 +154,7 @@ steps:
  - pytest -v -s entrypoints/multimodal

 - label: Entrypoints Integration (Pooling)
+  device: h200_35gb
  key: entrypoints-integration-pooling
  timeout_in_minutes: 50
  working_dir: "/vllm-workspace/tests"
@@ -15,6 +15,7 @@ steps:
    - pytest -v -s tests/kernels/ir

 - label: Kernels Core Operation Test
+  device: h200_35gb
  key: kernels-core-operation-test
  timeout_in_minutes: 120
  source_file_dependencies:
@@ -163,6 +164,7 @@ steps:
      - image-build-amd

 - label: Kernels Mamba Test
+  device: h200_35gb
  key: kernels-mamba-test
  timeout_in_minutes: 40
  source_file_dependencies:
@@ -235,6 +237,11 @@ steps:
  - vllm/model_executor/kernels/linear/cute_dsl/ll_bf16.py
  - vllm/model_executor/kernels/linear/cute_dsl/_ll_bf16_dotprod.py
  - vllm/model_executor/kernels/linear/cute_dsl/_ll_bf16_splitk.py
+  - vllm/cute_utils/
+  - vllm/model_executor/layers/mamba/ops/gdn_chunk_cutedsl/
+  - vllm/model_executor/layers/fused_moe/router/bf16x3_router_gemm_cutedsl.py
+  - tests/kernels/mamba/test_gdn_prefill_cutedsl.py
+  - tests/kernels/test_bf16x3_router_gemm_cutedsl.py
  - tests/kernels/test_ll_bf16_gemm.py
  - tests/kernels/test_top_k_per_row.py
  commands:
@@ -264,6 +271,8 @@ steps:
    - pytest -v -s tests/kernels/moe/test_flashinfer_moe.py
    - pytest -v -s tests/kernels/moe/test_trtllm_nvfp4_moe.py
    - pytest -v -s tests/kernels/moe/test_cutedsl_moe.py
+    - pytest -v -s tests/kernels/mamba/test_gdn_prefill_cutedsl.py
+    - pytest -v -s tests/kernels/test_bf16x3_router_gemm_cutedsl.py
    - pytest -v -s tests/kernels/test_ll_bf16_gemm.py
    # e2e
    - pytest -v -s tests/models/quantization/test_nvfp4.py
@@ -78,6 +78,28 @@ steps:
  commands:
  - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-small-tp.txt

+- label: LM Eval PCP (4xB200)
+  key: lm-eval-pcp-4xb200
+  timeout_in_minutes: 360
+  device: b200-k8s
+  num_devices: 4
+  optional: true
+  source_file_dependencies:
+  - csrc/
+  - tests/evals/gsm8k/configs/GLM-5.2-NVFP4-TP2-PCP2-EP.yaml
+  - tests/evals/gsm8k/configs/GLM-5.2-NVFP4-TP1-PCP4-EP.yaml
+  - tests/evals/gsm8k/configs/models-pcp.txt
+  - vllm/model_executor/layers/quantization
+  - vllm/config/parallel.py
+  - vllm/distributed/parallel_state.py
+  - vllm/model_executor/layers/attention/mla_attention.py
+  - vllm/model_executor/layers/attention/pcp.py
+  - vllm/v1/worker/gpu/model_runner.py
+  - vllm/v1/worker/gpu/pcp_manager.py
+  autorun_on_main: true
+  commands:
+  - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-pcp.txt
+
 - label: LM Eval Large Models EP (2xB200)
  key: lm-eval-large-models-ep-2xb200
  timeout_in_minutes: 60
@@ -64,8 +64,9 @@ steps:
      - image-build-amd

 - label: V1 Core + KV + Metrics
+  device: h200_35gb
  key: v1-core-kv-metrics
-  timeout_in_minutes: 60
+  timeout_in_minutes: 80
  source_file_dependencies:
    - vllm/config/
    - vllm/distributed/
@@ -3,6 +3,7 @@ depends_on:
  - image-build
 steps:
 - label: Model Executor
+  device: h200_35gb
  key: model-executor
  timeout_in_minutes: 45
  source_file_dependencies:
@@ -21,6 +21,7 @@ steps:
      - image-build-amd

 - label: Language Models Tests (Extra Standard) %N
+  device: h200_35gb
  key: language-models-tests-extra-standard
  timeout_in_minutes: 40
  source_file_dependencies:
@@ -51,8 +52,8 @@ steps:
      - tests/models/language/pooling/test_classification.py
      - vllm/_aiter_ops.py
      - vllm/platforms/rocm.py
-
 - label: Language Models Tests (Hybrid) %N
+  device: h200_35gb
  key: language-models-tests-hybrid
  timeout_in_minutes: 65
  source_file_dependencies:
@@ -63,8 +64,8 @@ steps:
    # Note: also needed to run plamo2 model in vLLM
    - uv pip install --system --no-build-isolation 'git+https://github.com/state-spaces/mamba@v2.3.0'
    - uv pip install --system --no-build-isolation 'git+https://github.com/Dao-AILab/causal-conv1d@v1.6.0'
-    # Shard hybrid language model tests
-    - pytest -v -s models/language/generation -m hybrid_model --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
+    # Shard the hybrid language model tests that are numerically stable on Hopper.
+    - pytest -v -s models/language/generation -m hybrid_model -k 'not granite-4.0-tiny-preview' --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
  parallelism: 2
  mirror:
    amd:
@@ -77,6 +78,20 @@ steps:
      - uv pip install --system --no-build-isolation 'git+https://github.com/Dao-AILab/causal-conv1d@v1.6.0'
      - pytest -v -s models/language/generation -m hybrid_model --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB

+# Granite 4 hybrid generation is sensitive to hardware-specific Triton SSD
+# autotuning (https://github.com/vllm-project/vllm/issues/25194). Keep this one
+# correctness test on L4 until its H200 output matches the Transformers reference.
+- label: Language Models Tests (Granite L4 Compatibility)
+  key: language-models-tests-granite-l4-compatibility
+  timeout_in_minutes: 65
+  source_file_dependencies:
+  - vllm/
+  - tests/models/language/generation
+  commands:
+    - uv pip install --system --no-build-isolation 'git+https://github.com/state-spaces/mamba@v2.3.0'
+    - uv pip install --system --no-build-isolation 'git+https://github.com/Dao-AILab/causal-conv1d@v1.6.0'
+    - pytest -v -s models/language/generation -m hybrid_model -k 'granite-4.0-tiny-preview'
+
 - label: Language Models Test (Extended Generation) # 80min
  device: h200_35gb
  key: language-models-test-extended-generation
@@ -119,6 +119,7 @@ steps:
      - vllm/model_executor/model_loader/

 - label: Multi-Modal Models (Extended Generation 1)
+  device: h200_35gb
  key: multi-modal-models-extended-generation-1
  optional: true
  source_file_dependencies:
@@ -116,8 +116,9 @@ steps:
      - image-build-amd

 - label: PyTorch Fullgraph Smoke Test
+  device: h200_35gb
  key: pytorch-fullgraph-smoke-test
-  timeout_in_minutes: 60
+  timeout_in_minutes: 90
  source_file_dependencies:
  - vllm/__init__.py
  - vllm/_aiter_ops.py
@@ -149,7 +150,42 @@ steps:
  # as it is a heavy test that is covered in other steps.
  # Use `find` to launch multiple instances of pytest so that
  # they do not suffer from https://github.com/vllm-project/vllm/issues/28965
-  - "find compile/fullgraph/ -name 'test_*.py' -not -name 'test_full_graph.py' -print0 | xargs -0 -n1 -I{} pytest -s -v '{}'"
+  - "find compile/fullgraph/ -name 'test_*.py' -not -name 'test_full_cudagraph.py' -not -name 'test_full_graph.py' -print0 | xargs -0 -n1 -I{} pytest -s -v '{}'"
+
+# Hopper-only DeepSeek-V2-Lite cases in this file require two 29.3-GiB model
+# instances and cannot fit a 35GB MIG slice. L4 retains the original coverage:
+# those SM90 cases skip while the architecture-compatible cases still run.
+- label: PyTorch Fullgraph CUDAGraph (L4 Compatibility)
+  key: pytorch-fullgraph-cudagraph-l4-compatibility
+  timeout_in_minutes: 60
+  source_file_dependencies:
+  - vllm/__init__.py
+  - vllm/_aiter_ops.py
+  - vllm/_custom_ops.py
+  - vllm/compilation/
+  - vllm/config/
+  - vllm/distributed/
+  - vllm/engine/
+  - vllm/env_override.py
+  - vllm/envs.py
+  - vllm/forward_context.py
+  - vllm/inputs/
+  - vllm/ir/
+  - vllm/kernels/
+  - vllm/logger.py
+  - vllm/model_executor/
+  - vllm/multimodal/
+  - vllm/platforms/
+  - vllm/plugins/
+  - vllm/sampling_params.py
+  - vllm/sequence.py
+  - vllm/transformers_utils/
+  - vllm/triton_utils/
+  - vllm/utils/
+  - vllm/v1/
+  - tests/compile
+  commands:
+  - pytest -s -v compile/fullgraph/test_full_cudagraph.py

 - label: PyTorch Fullgraph
  key: pytorch-fullgraph
@@ -3,8 +3,11 @@ depends_on:
  - image-build
 steps:
 - label: Quantization
+  device: h200_35gb
  key: quantization
-  timeout_in_minutes: 60
+  timeout_in_minutes: 75
+  env:
+    VLLM_USE_V2_MODEL_RUNNER: "0"
  source_file_dependencies:
  - csrc/
  - vllm/model_executor/layers/quantization
@@ -19,9 +22,13 @@ steps:
  # TODO(jerryzh168): resolve the above comment
  - uv pip install --system torchao==0.17.0 --index-url https://download.pytorch.org/whl/cu130
  - uv pip install --system conch-triton-kernels
-  - VLLM_TEST_FORCE_LOAD_FORMAT=auto pytest -v -s quantization/ --ignore quantization/test_blackwell_moe.py
+  # The SM90-only checkpoint currently contains a removed weight_chan_scale
+  # parameter. It was not exercised by the previous L4 job.
+  - VLLM_TEST_FORCE_LOAD_FORMAT=auto pytest -v -s quantization/ --ignore quantization/test_blackwell_moe.py -k 'not test_compressed_tensors_w4a8_fp8' --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
+  parallelism: 8

 - label: Quantized Fusions
+  device: h200_35gb
  key: quantized-fusions
  timeout_in_minutes: 20
  source_file_dependencies:
@@ -52,10 +59,14 @@ steps:
    - pytest -s -v tests/quantization/test_blackwell_moe.py

 - label: Quantized Models Test
+  device: h200_35gb
  key: quantized-models-test
-  timeout_in_minutes: 50
+  timeout_in_minutes: 65
+  env:
+    VLLM_USE_V2_MODEL_RUNNER: "0"
  source_file_dependencies:
  - vllm/model_executor/layers/quantization
  - tests/models/quantization
  commands:
-    - pytest -v -s models/quantization
+    - pytest -v -s models/quantization --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
+  parallelism: 3
@@ -81,6 +81,7 @@ steps:
  - pytest -s entrypoints/openai/correctness/test_lmeval.py::test_lm_eval_accuracy_v1_engine

 - label: Rust Frontend Tool Use
+  device: h200_35gb
  timeout_in_minutes: 25
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
@@ -30,7 +30,7 @@ repos:
  - id: markdownlint-cli2
    language_version: lts
    args: [--fix]
-    exclude: ^CLAUDE\.md$
+    exclude: (^|/)CLAUDE\.md$
 - repo: https://github.com/rhysd/actionlint
  rev: v1.7.7
  hooks:
@@ -69,12 +69,11 @@ def make_inputs(total_tokens, num_reqs, block_size):
    # Output workspace
    dst = torch.zeros(total_tokens, HEAD_DIM, dtype=torch.bfloat16, device="cuda")

-    seq_lens_t = torch.tensor(seq_lens, dtype=torch.int32, device="cuda")
    workspace_starts_t = torch.tensor(
        workspace_starts, dtype=torch.int32, device="cuda"
    )

-    return cache, dst, block_table, seq_lens_t, workspace_starts_t
+    return cache, dst, block_table, workspace_starts_t


 def bench_scenario(label, num_reqs, total_tokens_list, save_path):
@@ -94,7 +93,7 @@ def bench_scenario(label, num_reqs, total_tokens_list, save_path):
        )
    )
    def bench_fn(total_tokens, provider, num_reqs):
-        cache, dst, block_table, seq_lens_t, ws_starts = make_inputs(
+        cache, dst, block_table, ws_starts = make_inputs(
            total_tokens, num_reqs, BLOCK_SIZE
        )

@@ -102,7 +101,7 @@ def bench_scenario(label, num_reqs, total_tokens_list, save_path):

        ms, min_ms, max_ms = triton.testing.do_bench_cudagraph(
            lambda: ops.cp_gather_and_upconvert_fp8_kv_cache(
-                cache, dst, block_table, seq_lens_t, ws_starts, num_reqs
+                cache, dst, block_table, ws_starts, num_reqs
            ),
            quantiles=quantiles,
            rep=500,
@@ -430,6 +430,7 @@ set(VLLM_EXT_SRC
    "csrc/cpu/layernorm.cpp"
    "csrc/cpu/mla_decode.cpp"
    "csrc/cpu/pos_encoding.cpp"
+    "csrc/cpu/mamba_cpu.cpp"
    "csrc/moe/dynamic_4bit_int_moe_cpu.cpp"
    "csrc/cpu/cpu_attn.cpp"
    "csrc/cpu/torch_bindings.cpp")
@@ -489,6 +490,7 @@ if (ENABLE_X86_ISA)
        "csrc/cpu/spec_decode_utils.cpp"
        "csrc/cpu/cpu_attn.cpp"
        "csrc/cpu/dnnl_kernels.cpp"
+        "csrc/cpu/mamba_cpu.cpp"
        "csrc/cpu/torch_bindings.cpp"
        # TODO: Remove these files
        "csrc/cpu/activation.cpp"
@@ -502,6 +504,7 @@ if (ENABLE_X86_ISA)
        "csrc/cpu/utils.cpp"
        "csrc/cpu/spec_decode_utils.cpp"
        "csrc/cpu/cpu_attn.cpp"
+        "csrc/cpu/mamba_cpu.cpp"
        "csrc/cpu/dnnl_kernels.cpp"
        "csrc/cpu/torch_bindings.cpp"
        # TODO: Remove these files
@@ -14,7 +14,7 @@ else()
  FetchContent_Declare(
    tml_fa4
    GIT_REPOSITORY https://github.com/vllm-project/tml-fa4.git
-    GIT_TAG 13374f0c855acc1add1bf30444bd67aebbc24a8e
+    GIT_TAG b206834606ed5b5f21f8eed6b0683f528ea9cf7d
    GIT_PROGRESS TRUE
    CONFIGURE_COMMAND ""
    BUILD_COMMAND "")
@@ -67,9 +67,8 @@ void cp_gather_and_upconvert_fp8_kv_cache(
    torch::Tensor const& src_cache,         // [NUM_BLOCKS, BLOCK_SIZE, 656]
    torch::Tensor const& dst,               // [TOT_TOKENS, 576]
    torch::Tensor const& block_table,       // [BATCH, BLOCK_INDICES]
-    torch::Tensor const& seq_lens,          // [BATCH]
    torch::Tensor const& workspace_starts,  // [BATCH]
-    int64_t batch_size);
+    int64_t batch_size, std::optional<torch::Tensor> seq_starts = std::nullopt);

 // Indexer K quantization and cache function
 void indexer_k_quant_and_cache(
@@ -336,13 +336,14 @@ struct FP32Vec8 : public Vec<FP32Vec8> {
    reg.val[1] = fp16_to_fp32_bits(raw_lo);
  }
  float reduce_sum() const {
-    AliasReg ar;
-    ar.reg = reg;
-    float result = 0;
-    unroll_loop<int, VEC_ELEM_NUM>(
-        [&result, &ar](int i) { result += ar.values[i]; });
-
-    return result;
+    // VSX horizontal reduction: 3 vector ops instead of 8 scalar adds.
+    // Step 1: pairwise sum of the two 4-wide halves
+    __vector float s = vec_add(reg.val[0], reg.val[1]);
+    // Step 2: rotate by 8 bytes (2 floats) and add
+    s = vec_add(s, vec_sld(s, s, 8));
+    // Step 3: rotate by 4 bytes (1 float) and add  => all lanes hold total
+    s = vec_add(s, vec_sld(s, s, 4));
+    return vec_extract(s, 0);
  }
  FP32Vec8 exp() const {
    f32x4x2_t out;
@@ -0,0 +1,285 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+//
+// CPU at::Tensor wrappers for Mamba decode-step kernels defined in
+// mamba_kernels.hpp.
+
+#include "cpu/mamba_kernels.hpp"
+
+#include <ATen/ATen.h>
+#include <torch/library.h>
+#include <c10/util/Optional.h>
+
+#include "cpu_types.hpp"
+
+// ---------------------------------------------------------------------------
+// causal_conv1d_update
+// ---------------------------------------------------------------------------
+at::Tensor causal_conv1d_update_cpu_impl(
+    at::Tensor& x, at::Tensor& conv_state, const at::Tensor& weight,
+    const c10::optional<at::Tensor>& bias,
+    const c10::optional<std::string>& activation,
+    const c10::optional<at::Tensor>& conv_state_indices,
+    const c10::optional<at::Tensor>& query_start_loc, int64_t pad_slot_id) {
+  bool do_silu = false;
+  if (activation.has_value()) {
+    const std::string& act = activation.value();
+    do_silu = (act == "silu" || act == "swish");
+  }
+
+  at::ScalarType dtype = x.scalar_type();
+
+  // Input x: contiguous in native dtype.
+  at::Tensor x_c = x.is_contiguous() ? x : x.contiguous();
+
+  // conv_state: NEVER copy the full paged tensor just for layout reasons.
+  // If the dtype matches we work directly on conv_state (contiguous or not)
+  // by extracting strides and passing them to the kernel.
+  // Only a dtype-conversion copy is made when types differ (rare for BF16).
+  bool state_type_ok = (conv_state.scalar_type() == dtype);
+  at::Tensor state_c = state_type_ok ? conv_state : conv_state.to(dtype);
+  // state_c and conv_state may be non-contiguous — that is intentional.
+
+  // Weight: coerce to same dtype if needed (should match in practice)
+  at::Tensor w_c =
+      (weight.scalar_type() != dtype)
+          ? weight.to(dtype).contiguous()
+          : (weight.is_contiguous() ? weight : weight.contiguous());
+
+  // Bias stays float32 (small scalar, used only for fp32 accumulation)
+  at::Tensor bias_f32;
+  if (bias.has_value() && bias.value().defined())
+    bias_f32 = bias.value().to(at::kFloat).contiguous();
+
+  int64_t batch = x_c.size(0);
+  int64_t dim = x_c.size(1);
+  int64_t seqlen = (x_c.dim() == 3) ? x_c.size(2) : 1;
+  int64_t width = w_c.size(1);
+  int64_t state_len = state_c.size(2);
+
+  // Extract strides — works for contiguous AND non-contiguous (transposed)
+  // state. stride(0): between cache slots (e.g. num_slots × dim × width-1 in
+  // contiguous) stride(1): between conv channels (dim stride) stride(2):
+  // between state elements (=1 when contiguous, =dim when transposed)
+  int64_t stride_s_slot = state_c.stride(0);
+  int64_t stride_s_dim = state_c.stride(1);
+  int64_t stride_s_state = state_c.stride(2);
+
+  at::Tensor out = x_c.clone();  // native dtype, no float32 alloc
+
+  const int32_t* cache_idx_ptr = nullptr;
+  at::Tensor cache_idx_int;
+  if (conv_state_indices.has_value()) {
+    cache_idx_int = conv_state_indices.value().to(at::kInt).contiguous();
+    cache_idx_ptr = cache_idx_int.data_ptr<int32_t>();
+  }
+
+  VLLM_DISPATCH_FLOATING_TYPES(dtype, "causal_conv1d_update", [&] {
+    mamba_cpu::causal_conv1d_update_kernel<scalar_t>(
+        x_c.data_ptr<scalar_t>(), state_c.data_ptr<scalar_t>(), stride_s_slot,
+        stride_s_dim, stride_s_state, w_c.data_ptr<scalar_t>(),
+        bias_f32.defined() ? bias_f32.data_ptr<float>() : nullptr,
+        out.data_ptr<scalar_t>(), cache_idx_ptr,
+        static_cast<int32_t>(pad_slot_id), batch, dim, seqlen, width, state_len,
+        do_silu);
+  });
+
+  // Write back only when a type-conversion copy was made.
+  // Layout-only non-contiguity is handled via strides above — no copy needed.
+  if (!state_type_ok) conv_state.copy_(state_c);
+
+  return out;
+}
+
+// ---------------------------------------------------------------------------
+// selective_state_update
+// ---------------------------------------------------------------------------
+void selective_state_update_cpu_impl(
+    at::Tensor& state,    // (nstates, nheads, dim, dstate)
+    const at::Tensor& x,  // (N, nheads, dim)
+    const at::Tensor& dt, const at::Tensor& A, const at::Tensor& B,
+    const at::Tensor& C, const c10::optional<at::Tensor>& D,
+    const c10::optional<at::Tensor>& z,
+    const c10::optional<at::Tensor>& dt_bias, bool dt_softplus,
+    const c10::optional<at::Tensor>& state_batch_indices,
+    const c10::optional<at::Tensor>& dst_state_batch_indices,
+    int64_t null_block_id, at::Tensor& out,
+    const c10::optional<at::Tensor>& num_accepted_tokens,
+    const c10::optional<at::Tensor>& cu_seqlens) {
+  at::ScalarType state_type = state.scalar_type();
+  at::ScalarType input_type = x.scalar_type();
+
+  // x, B, C must be contiguous and match input_type
+  auto ensure_input = [input_type](const at::Tensor& t) -> at::Tensor {
+    at::Tensor r = (t.scalar_type() != input_type) ? t.to(input_type) : t;
+    return r.is_contiguous() ? r : r.contiguous();
+  };
+  at::Tensor x_in = ensure_input(x);
+  at::Tensor B_in = ensure_input(B);
+  at::Tensor C_in = ensure_input(C);
+  at::Tensor z_in;
+  if (z.has_value() && z.value().defined()) z_in = ensure_input(z.value());
+
+  // A, D, dt_bias are float32 model parameters that arrive here as expanded
+  // tensors, e.g. A is (nheads, head_dim, dstate) with strides (1, 0, 0).
+  // We need just the scalar value per head as a (nheads,) 1-D array so that
+  // A_ptr[h] in the kernel correctly reads head h's value.
+  //
+  // Strategy: peel trailing expanded (stride=0) dims via .select(), which is
+  // a zero-copy view.  For A: (nheads, head_dim, dstate) strides (1,0,0)
+  //   → .select(2,0) → (nheads, head_dim) strides (1,0)
+  //   → .select(1,0) → (nheads,)          stride  (1,)   ← contiguous, free.
+  // No allocation, no type conversion (A is already float32).
+  auto to_per_head_1d_f32 = [](const at::Tensor& t) -> at::Tensor {
+    at::Tensor r = t;
+    // Peel trailing dimensions that are broadcast (stride=0 or size=1)
+    while (r.dim() > 1) r = r.select(r.dim() - 1, 0);
+    if (r.scalar_type() != at::kFloat) r = r.to(at::kFloat);
+    return r.is_contiguous() ? r : r.contiguous();
+  };
+
+  at::Tensor A_f32 = to_per_head_1d_f32(A);  // (nheads,) float32
+  at::Tensor D_f32, dt_bias_f32;
+  if (D.has_value() && D.value().defined())
+    D_f32 = to_per_head_1d_f32(D.value());
+  if (dt_bias.has_value() && dt_bias.value().defined())
+    dt_bias_f32 = to_per_head_1d_f32(dt_bias.value());
+
+  // dt: reduce (N, nheads, head_dim) expanded tensor → (N, nheads) BEFORE
+  // the type conversion so we convert head_dim x fewer elements.
+  at::Tensor dt_f32;
+  {
+    // If dt was expanded to (N, nheads, head_dim) with stride-0 in dim 2,
+    // take a zero-copy view of index 0 along that dim first.
+    at::Tensor t2 = (dt.dim() == 3) ? dt.select(2, 0) : dt;  // (N, nheads)
+    at::Tensor t3 = (t2.scalar_type() != at::kFloat) ? t2.to(at::kFloat) : t2;
+    dt_f32 = t3.is_contiguous() ? t3 : t3.contiguous();
+  }
+
+  int64_t nheads = state.size(1);
+  int64_t dim = state.size(2);
+  int64_t dstate = state.size(3);
+  int64_t N = (cu_seqlens.has_value() && cu_seqlens.value().defined())
+                  ? cu_seqlens.value().size(0) - 1
+                  : x_in.size(0);
+  int64_t ngroups = B_in.size(1);
+
+  // Strides
+  int64_t stride_state_n = state.stride(0);
+  int64_t stride_state_h = state.stride(1);
+  int64_t stride_state_d = state.stride(2);
+  int64_t stride_x_n = x_in.stride(0);
+  int64_t stride_x_h = x_in.stride(1);
+  int64_t stride_dt_n = dt_f32.stride(0);  // dt is (N, nheads)
+  int64_t stride_BC_n = B_in.stride(0);
+  int64_t stride_BC_g = B_in.stride(1);
+  int64_t stride_out_n = out.stride(0);
+  int64_t stride_out_h = out.stride(1);
+
+  // Optional index pointers
+  auto get_int32_ptr =
+      [](const c10::optional<at::Tensor>& opt) -> const int32_t* {
+    return (opt.has_value() && opt.value().defined())
+               ? opt.value().data_ptr<int32_t>()
+               : nullptr;
+  };
+  const int32_t* sbi_ptr = get_int32_ptr(state_batch_indices);
+  const int32_t* dsbi_ptr = get_int32_ptr(dst_state_batch_indices);
+  const int32_t* nat_ptr = get_int32_ptr(num_accepted_tokens);
+  const int32_t* csl_ptr = get_int32_ptr(cu_seqlens);
+
+  // Dispatch on (state_t, input_t, out_t): write directly into `out`
+  // without any intermediate float32 buffer.
+  VLLM_DISPATCH_FLOATING_TYPES(state_type, "ssu_state", [&] {
+    using state_t = scalar_t;
+    VLLM_DISPATCH_FLOATING_TYPES(input_type, "ssu_input", [&] {
+      using input_t = scalar_t;
+      VLLM_DISPATCH_FLOATING_TYPES(out.scalar_type(), "ssu_out", [&] {
+        using out_t = scalar_t;
+        mamba_cpu::selective_state_update_kernel<state_t, input_t, out_t>(
+            state.data_ptr<state_t>(), stride_state_n, stride_state_h,
+            stride_state_d, x_in.data_ptr<input_t>(), stride_x_n, stride_x_h,
+            dt_f32.data_ptr<float>(), stride_dt_n, A_f32.data_ptr<float>(),
+            B_in.data_ptr<input_t>(), C_in.data_ptr<input_t>(), stride_BC_n,
+            stride_BC_g, D_f32.defined() ? D_f32.data_ptr<float>() : nullptr,
+            z_in.defined() ? z_in.data_ptr<input_t>() : nullptr,
+            dt_bias_f32.defined() ? dt_bias_f32.data_ptr<float>() : nullptr,
+            out.data_ptr<out_t>(), stride_out_n, stride_out_h, sbi_ptr,
+            dsbi_ptr, static_cast<int32_t>(null_block_id), nat_ptr, csl_ptr, N,
+            nheads, ngroups, dim, dstate, dt_softplus);
+      });
+    });
+  });
+}
+
+// ---------------------------------------------------------------------------
+// mamba_chunk_scan_fwd_cpu
+// ---------------------------------------------------------------------------
+void mamba_chunk_scan_fwd_cpu_impl(
+    at::Tensor& out,  // [seqlen, nheads, headdim] — pre-allocated by caller
+    at::Tensor&
+        final_states,     // [batch, nheads, headdim, dstate] float32 contiguous
+    const at::Tensor& x,  // [seqlen, nheads, headdim]
+    const at::Tensor&
+        dt,  // [seqlen, nheads] float32 (preprocessed: bias+softplus+clamp)
+    const at::Tensor& A,                 // [nheads] float32
+    const at::Tensor& B,                 // [seqlen, ngroups, dstate]
+    const at::Tensor& C,                 // [seqlen, ngroups, dstate]
+    const c10::optional<at::Tensor>& D,  // [nheads] float32 (optional)
+    const c10::optional<at::Tensor>& z,  // [seqlen, nheads, headdim] (optional)
+    const at::Tensor& cu_seqlens         // [batch+1] int32
+) {
+  const at::ScalarType input_type = x.scalar_type();
+
+  auto ensure_contig = [input_type](const at::Tensor& t) -> at::Tensor {
+    at::Tensor r = (t.scalar_type() != input_type) ? t.to(input_type) : t;
+    return r.is_contiguous() ? r : r.contiguous();
+  };
+  at::Tensor x_in = ensure_contig(x);
+  at::Tensor B_in = ensure_contig(B);
+  at::Tensor C_in = ensure_contig(C);
+  at::Tensor z_in;
+  if (z.has_value() && z.value().defined()) z_in = ensure_contig(z.value());
+
+  // A and D are float32 model parameters, potentially broadcast-expanded.
+  // Strip trailing broadcast dims to get a contiguous (nheads,) array.
+  auto to_per_head_f32 = [](const at::Tensor& t) -> at::Tensor {
+    at::Tensor r = t;
+    while (r.dim() > 1) r = r.select(r.dim() - 1, 0);
+    if (r.scalar_type() != at::kFloat) r = r.to(at::kFloat);
+    return r.is_contiguous() ? r : r.contiguous();
+  };
+  at::Tensor A_f32 = to_per_head_f32(A);
+  at::Tensor D_f32;
+  if (D.has_value() && D.value().defined()) D_f32 = to_per_head_f32(D.value());
+
+  // dt: [seqlen, nheads] float32 — caller has applied bias+softplus+clamp in
+  // Python.
+  at::Tensor dt_c = dt.is_contiguous() ? dt : dt.contiguous();
+  if (dt_c.scalar_type() != at::kFloat) dt_c = dt_c.to(at::kFloat);
+
+  at::Tensor cu_int = cu_seqlens.to(at::kInt).contiguous();
+
+  const int64_t batch = final_states.size(0);
+  const int64_t nheads = final_states.size(1);
+  const int64_t headdim = final_states.size(2);
+  const int64_t dstate = final_states.size(3);
+  const int64_t ngroups = B_in.size(1);
+
+  TORCH_CHECK(final_states.is_contiguous(),
+              "mamba_chunk_scan_fwd_cpu: final_states must be contiguous");
+  TORCH_CHECK(out.is_contiguous(),
+              "mamba_chunk_scan_fwd_cpu: out must be contiguous (writes via "
+              "raw data_ptr)");
+
+  VLLM_DISPATCH_FLOATING_TYPES(input_type, "mamba_chunk_scan_fwd_cpu", [&] {
+    mamba_cpu::mamba_chunk_scan_fwd_kernel<scalar_t>(
+        final_states.data_ptr<float>(), x_in.data_ptr<scalar_t>(),
+        dt_c.data_ptr<float>(), A_f32.data_ptr<float>(),
+        B_in.data_ptr<scalar_t>(), C_in.data_ptr<scalar_t>(),
+        D_f32.defined() ? D_f32.data_ptr<float>() : nullptr,
+        z_in.defined() ? z_in.data_ptr<scalar_t>() : nullptr,
+        out.data_ptr<scalar_t>(), cu_int.data_ptr<int32_t>(), batch, nheads,
+        ngroups, headdim, dstate);
+  });
+}
@@ -0,0 +1,382 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+//
+// Fused CPU vector kernels for Mamba decode-step hotspots:
+//   - causal_conv1d_update  (depthwise 1-D conv state roll + compute)
+//   - selective_state_update (SSM recurrence, single-step)
+
+#pragma once
+
+#include "cpu_types.hpp"
+#include <cmath>
+#include <cstring>
+#include <cstdint>
+#include <algorithm>
+
+namespace mamba_cpu {
+
+// ---------------------------------------------------------------------------
+// causal_conv1d_update — templated for native BF16/FP32
+//
+// state_ptr may point to a NON-CONTIGUOUS paged KV cache tensor.
+// Explicit strides are passed so the kernel writes directly into the
+// correct memory locations without making a contiguous copy of the full
+// paged tensor (which was the source of the 34-41% direct_copy_kernel).
+//
+//   stride_s_slot  = state.stride(0)  — between cache slots
+//   stride_s_dim   = state.stride(1)  — between conv_dim channels
+//   stride_s_state = state.stride(2)  — between state elements
+//
+// When stride_s_state == 1 (contiguous), the memmove fast path is used.
+// ---------------------------------------------------------------------------
+template <typename scalar_t>
+inline void causal_conv1d_update_kernel(
+    const scalar_t* __restrict__ x_ptr, scalar_t* __restrict__ state_ptr,
+    int64_t stride_s_slot, int64_t stride_s_dim, int64_t stride_s_state,
+    const scalar_t* __restrict__ weight_ptr, const float* __restrict__ bias_ptr,
+    scalar_t* __restrict__ out_ptr, const int32_t* __restrict__ cache_idxs,
+    int32_t pad_slot_id, int64_t batch, int64_t dim, int64_t seqlen,
+    int64_t width, int64_t state_len, bool do_silu) {
+#pragma omp parallel for
+  for (int64_t b = 0; b < batch; ++b) {
+    int64_t cache_idx = (cache_idxs != nullptr) ? cache_idxs[b] : b;
+    if (cache_idx == pad_slot_id) continue;
+
+    for (int64_t t = 0; t < seqlen; ++t) {
+      const scalar_t* x_b = x_ptr + (b * dim * seqlen + t);
+      scalar_t* out_b = out_ptr + (b * dim * seqlen + t);
+      // Base of this slot in the (possibly non-contiguous) paged state
+      scalar_t* s_base = state_ptr + cache_idx * stride_s_slot;
+
+      for (int64_t d = 0; d < dim; ++d) {
+        float x_val = static_cast<float>(x_b[d * seqlen]);
+        scalar_t* sd = s_base + d * stride_s_dim;  // start of this dim's state
+        const scalar_t* w = weight_ptr + d * width;
+
+        // Accumulate in float32 for precision
+        float acc = (bias_ptr != nullptr) ? bias_ptr[d] : 0.0f;
+        for (int64_t k = 0; k < state_len; ++k) {
+          acc += static_cast<float>(w[k]) *
+                 static_cast<float>(sd[k * stride_s_state]);
+        }
+        acc += static_cast<float>(w[state_len]) * x_val;
+
+        // Shift state left and append new input.
+        // Use memmove when contiguous (stride==1); element loop otherwise.
+        if (stride_s_state == 1) {
+          if (state_len > 1)
+            std::memmove(sd, sd + 1, (state_len - 1) * sizeof(scalar_t));
+          if (state_len > 0) sd[state_len - 1] = static_cast<scalar_t>(x_val);
+        } else {
+          for (int64_t k = 0; k < state_len - 1; ++k)
+            sd[k * stride_s_state] = sd[(k + 1) * stride_s_state];
+          if (state_len > 0)
+            sd[(state_len - 1) * stride_s_state] = static_cast<scalar_t>(x_val);
+        }
+
+        if (do_silu) {
+          float sigmoid = (acc >= 0) ? 1.0f / (1.0f + std::exp(-acc))
+                                     : std::exp(acc) / (1.0f + std::exp(acc));
+          acc *= sigmoid;
+        }
+        out_b[d * seqlen] = static_cast<scalar_t>(acc);
+      }
+    }
+  }
+}
+
+// ---------------------------------------------------------------------------
+// selective_state_update
+//
+// Template parameters:
+//   state_t  - dtype of ssm_state cache (typically BFloat16)
+//   input_t  - dtype of x, B, C (typically BFloat16)
+//   out_t    - dtype of output tensor (typically BFloat16)
+//              Write directly — no float32 intermediate buffer needed.
+//
+// A, D, dt_bias are accepted as const float* (they are always float32
+// model parameters in Mamba2).  This eliminates the per-call float32→BF16
+// conversion and the .contiguous() materialisation of the broadcast-expand.
+//
+// dt is accepted as a (N, nheads) scalar-per-head tensor, not as the
+// (N, nheads, head_dim) expansion, so no .contiguous() copy is needed.
+// ---------------------------------------------------------------------------
+template <typename state_t, typename input_t, typename out_t = float>
+inline void selective_state_update_kernel(
+    state_t* __restrict__ state_ptr, int64_t stride_state_n,
+    int64_t stride_state_h, int64_t stride_state_d,
+    const input_t* __restrict__ x_ptr, int64_t stride_x_n, int64_t stride_x_h,
+    // dt: (N, nheads) — scalar per head, NOT expanded to head_dim
+    const float* __restrict__ dt_ptr, int64_t stride_dt_n,
+    // A: (nheads,) float32 — scalar per head
+    const float* __restrict__ A_ptr, const input_t* __restrict__ B_ptr,
+    const input_t* __restrict__ C_ptr, int64_t stride_BC_n, int64_t stride_BC_g,
+    // D: (nheads,) float32 — scalar per head (nullptr if not used)
+    const float* __restrict__ D_ptr,
+    // z: same shape as x (optional)
+    const input_t* __restrict__ z_ptr,
+    // dt_bias: (nheads,) float32 — scalar per head (nullptr if not used)
+    const float* __restrict__ dt_bias_ptr, out_t* __restrict__ out_ptr,
+    int64_t stride_out_n, int64_t stride_out_h,
+    const int32_t* __restrict__ state_batch_indices,
+    const int32_t* __restrict__ dst_state_batch_indices, int32_t null_block_id,
+    const int32_t* __restrict__ num_accepted_tokens,
+    const int32_t* __restrict__ cu_seqlens, int64_t N, int64_t nheads,
+    int64_t ngroups, int64_t dim, int64_t dstate, bool dt_softplus) {
+  using state_vec_t = vec_op::vec_t<state_t>;
+  using input_vec_t = vec_op::vec_t<input_t>;
+  constexpr int VEC_ELEM_NUM = 8;
+
+  int64_t nheads_per_group = nheads / ngroups;
+
+  for (int64_t seq_idx = 0; seq_idx < N; ++seq_idx) {
+    int64_t bos, seq_len;
+    if (cu_seqlens != nullptr) {
+      bos = cu_seqlens[seq_idx];
+      seq_len = cu_seqlens[seq_idx + 1] - bos;
+    } else {
+      bos = seq_idx;
+      seq_len = 1;
+    }
+
+    int64_t state_read_idx = (state_batch_indices != nullptr)
+                                 ? state_batch_indices[seq_idx]
+                                 : seq_idx;
+    if (state_read_idx == null_block_id) continue;
+
+    int64_t state_write_idx = (num_accepted_tokens == nullptr)
+                                  ? ((dst_state_batch_indices != nullptr)
+                                         ? dst_state_batch_indices[seq_idx]
+                                         : state_read_idx)
+                                  : -1;
+
+    state_t* s = state_ptr + state_read_idx * stride_state_n;
+
+    for (int64_t t = 0; t < seq_len; ++t) {
+      int64_t token_idx = bos + t;
+      const input_t* x_tok = x_ptr + token_idx * stride_x_n;
+      // dt: (N, nheads) — one float per head per token
+      const float* dt_tok = dt_ptr + token_idx * stride_dt_n;
+      const input_t* B_tok = B_ptr + token_idx * stride_BC_n;
+      const input_t* C_tok = C_ptr + token_idx * stride_BC_n;
+      out_t* out_tok = out_ptr + token_idx * stride_out_n;
+
+#pragma omp parallel for
+      for (int64_t h = 0; h < nheads; ++h) {
+        int64_t g = h / nheads_per_group;
+        const input_t* x_h = x_tok + h * stride_x_h;
+        const input_t* B_g = B_tok + g * stride_BC_g;
+        const input_t* C_g = C_tok + g * stride_BC_g;
+        out_t* out_h = out_tok + h * stride_out_h;
+        state_t* s_h = s + h * stride_state_h;
+
+        // Read scalars-per-head (A, dt, dt_bias, D) — no per-dim indexing
+        float dt_val = dt_tok[h];
+        if (dt_bias_ptr != nullptr) dt_val += dt_bias_ptr[h];
+        if (dt_softplus) {
+          dt_val = (dt_val <= 20.0f) ? std::log1p(std::exp(dt_val)) : dt_val;
+        }
+        const float A_val = A_ptr[h];  // scalar: same for all dim, dstate
+        const float D_val = (D_ptr != nullptr) ? D_ptr[h] : 0.0f;
+
+        const input_t* z_h =
+            (z_ptr != nullptr) ? z_ptr + token_idx * stride_x_n + h * stride_x_h
+                               : nullptr;
+
+        vec_op::FP32Vec8 dt_vec(dt_val);
+        // dA = exp(A * dt): A and dt are SCALARS per head, so compute once
+        // and broadcast.  This saves 7 redundant std::exp() calls that
+        // FP32Vec8::exp() would otherwise make on the broadcast vector.
+        const float dA_scalar = std::exp(A_val * dt_val);
+        vec_op::FP32Vec8 dA(dA_scalar);  // broadcast
+
+        for (int64_t d = 0; d < dim; ++d) {
+          float x_val = static_cast<float>(x_h[d]);
+
+          vec_op::FP32Vec8 out_vec(0.0f);
+          state_t* s_hd = s_h + d * stride_state_d;
+          const input_t* B_g_base = B_g;
+          const input_t* C_g_base = C_g;
+
+          vec_op::FP32Vec8 x_vec(x_val);
+          // dBx = B * x * dt — same dA for all dstate (A is scalar)
+          // s_new = s * dA + B * x * dt
+
+          int64_t n = 0;
+          for (; n <= dstate - VEC_ELEM_NUM; n += VEC_ELEM_NUM) {
+            vec_op::FP32Vec8 B_v((input_vec_t(B_g_base + n)));
+            vec_op::FP32Vec8 C_v((input_vec_t(C_g_base + n)));
+            vec_op::FP32Vec8 s_v((state_vec_t(s_hd + n)));
+
+            vec_op::FP32Vec8 dBx = B_v * x_vec * dt_vec;
+            vec_op::FP32Vec8 s_new = s_v * dA + dBx;
+
+            state_vec_t(s_new).save(s_hd + n);
+            out_vec = out_vec + s_new * C_v;
+          }
+
+          float out_val = out_vec.reduce_sum();
+          for (; n < dstate; ++n) {
+            // Reuse dA_scalar computed once per head — no exp() re-call
+            float dBx = static_cast<float>(B_g[n]) * x_val * dt_val;
+            float s_new = static_cast<float>(s_hd[n]) * dA_scalar + dBx;
+            s_hd[n] = static_cast<state_t>(s_new);
+            out_val += s_new * static_cast<float>(C_g[n]);
+          }
+
+          if (D_ptr != nullptr) out_val += x_val * D_val;
+          if (z_h != nullptr) {
+            float z_val = static_cast<float>(z_h[d]);
+            float sigmoid = (z_val >= 0)
+                                ? 1.0f / (1.0f + std::exp(-z_val))
+                                : std::exp(z_val) / (1.0f + std::exp(z_val));
+            out_val *= z_val * sigmoid;
+          }
+          out_h[d] = static_cast<out_t>(out_val);
+        }
+      }
+
+      if (num_accepted_tokens != nullptr &&
+          dst_state_batch_indices != nullptr) {
+        int64_t token_dst_idx = dst_state_batch_indices[seq_idx * seq_len + t];
+        if (token_dst_idx != null_block_id && token_dst_idx != state_read_idx) {
+          state_t* dst_s = state_ptr + token_dst_idx * stride_state_n;
+          std::memmove(dst_s, s, nheads * stride_state_h * sizeof(state_t));
+        }
+      }
+    }
+
+    if (num_accepted_tokens == nullptr && state_write_idx != null_block_id &&
+        state_write_idx != state_read_idx) {
+      state_t* dst_s = state_ptr + state_write_idx * stride_state_n;
+      std::memmove(dst_s, s, nheads * stride_state_h * sizeof(state_t));
+    }
+  }
+}
+
+// ---------------------------------------------------------------------------
+// mamba_chunk_scan_fwd
+//
+// Prefill SSM recurrence for Mamba2 / SSD models.
+//
+// Key difference from selective_state_update_kernel (decode path):
+//   - #pragma omp parallel for collapse(2) is OUTSIDE the time loop.
+//     Each thread owns a (batch, head) slice and runs the entire token
+//     sequence without any per-token OpenMP synchronisation overhead.
+//     For seqlen=256, this eliminates 256 thread-barrier launches per batch.
+//
+// `dt` arrives already processed (float32, after bias + softplus + clamp)
+// to keep this kernel simple. Preprocessing is done in the Python wrapper.
+//
+// `states_ptr` points to the [batch, nheads, headdim, dstate] float32 output
+// tensor, pre-initialised by the caller (zero or from initial_states).
+// Each (b, h) slice is private to exactly one thread via collapse(2), so
+// there are no write conflicts.
+//
+// D is treated as a scalar per head ([nheads] float32).
+// ---------------------------------------------------------------------------
+template <typename input_t>
+inline void mamba_chunk_scan_fwd_kernel(
+    float* __restrict__ states_ptr,     // [batch, nheads, headdim, dstate] f32
+    const input_t* __restrict__ x_ptr,  // [seqlen, nheads, headdim]
+    const float* __restrict__ dt_ptr,   // [seqlen, nheads] f32 (preprocessed)
+    const float* __restrict__ A_ptr,    // [nheads] f32
+    const input_t* __restrict__ B_ptr,  // [seqlen, ngroups, dstate]
+    const input_t* __restrict__ C_ptr,  // [seqlen, ngroups, dstate]
+    const float* __restrict__ D_ptr,    // [nheads] f32 (nullable)
+    const input_t* __restrict__ z_ptr,  // [seqlen, nheads, headdim] (nullable)
+    input_t* __restrict__ out_ptr,      // [seqlen, nheads, headdim]
+    const int32_t* __restrict__ cu_seqlens,  // [batch+1] int32
+    int64_t batch, int64_t nheads, int64_t ngroups, int64_t headdim,
+    int64_t dstate) {
+  using input_vec_t = vec_op::vec_t<input_t>;
+  constexpr int VEC_ELEM_NUM = 8;
+
+  const int64_t nheads_per_group = nheads / ngroups;
+  // states layout: [batch, nheads, headdim, dstate] contiguous (caller
+  // guarantee)
+  const int64_t stride_s_b = nheads * headdim * dstate;
+  const int64_t stride_s_h = headdim * dstate;
+  // stride_s_d = dstate, stride_s_n = 1
+
+#pragma omp parallel for collapse(2) schedule(static)
+  for (int64_t b = 0; b < batch; ++b) {
+    for (int64_t h = 0; h < nheads; ++h) {
+      const int64_t seq_start = cu_seqlens[b];
+      const int64_t seq_end = cu_seqlens[b + 1];
+      const int64_t g = h / nheads_per_group;
+
+      const float A_val = A_ptr[h];
+      const float D_val = (D_ptr != nullptr) ? D_ptr[h] : 0.0f;
+
+      // Working state slice: states[b, h, :, :] — float32, headdim * dstate.
+      // Fits in L1/L2 for typical dims (e.g. 64*128*4 = 32 KB).
+      float* s_bh = states_ptr + b * stride_s_b + h * stride_s_h;
+
+      for (int64_t t = seq_start; t < seq_end; ++t) {
+        const input_t* x_h = x_ptr + t * nheads * headdim + h * headdim;
+        const float* dt_h = dt_ptr + t * nheads + h;
+        const input_t* B_g = B_ptr + t * ngroups * dstate + g * dstate;
+        const input_t* C_g = C_ptr + t * ngroups * dstate + g * dstate;
+        const input_t* z_h = (z_ptr != nullptr)
+                                 ? z_ptr + t * nheads * headdim + h * headdim
+                                 : nullptr;
+        input_t* out_h = out_ptr + t * nheads * headdim + h * headdim;
+
+        const float dt_val = *dt_h;
+        const float dA_val = std::exp(A_val * dt_val);
+        const vec_op::FP32Vec8 dA_vec(dA_val);  // broadcast scalar
+        const vec_op::FP32Vec8 dt_vec(dt_val);
+
+        for (int64_t d = 0; d < headdim; ++d) {
+          const float x_val = static_cast<float>(x_h[d]);
+          float* s_bhd = s_bh + d * dstate;  // [dstate] contiguous float32
+
+          // Vectorised SSM update + readout over dstate:
+          //   s_new = s * dA + x * dt * B
+          //   y    += s_new * C
+          int64_t n = 0;
+          vec_op::FP32Vec8 y_vec(0.0f);
+          const vec_op::FP32Vec8 x_vec(x_val);
+
+          for (; n <= dstate - VEC_ELEM_NUM; n += VEC_ELEM_NUM) {
+            const vec_op::FP32Vec8 B_v((input_vec_t(B_g + n)));
+            const vec_op::FP32Vec8 C_v((input_vec_t(C_g + n)));
+            const vec_op::FP32Vec8 s_v(s_bhd + n);
+
+            const vec_op::FP32Vec8 s_new = s_v * dA_vec + x_vec * dt_vec * B_v;
+            s_new.save(s_bhd + n);
+            y_vec = y_vec + s_new * C_v;
+          }
+
+          float y_val = y_vec.reduce_sum();
+
+          // Scalar tail for remaining dstate elements
+          for (; n < dstate; ++n) {
+            const float B_n = static_cast<float>(B_g[n]);
+            const float C_n = static_cast<float>(C_g[n]);
+            const float s_new = s_bhd[n] * dA_val + x_val * dt_val * B_n;
+            s_bhd[n] = s_new;
+            y_val += s_new * C_n;
+          }
+
+          // D skip connection (scalar per head)
+          if (D_ptr != nullptr) y_val += x_val * D_val;
+
+          // z gating: out = y * z * sigmoid(z)  (SiLU)
+          if (z_h != nullptr) {
+            const float z_val = static_cast<float>(z_h[d]);
+            const float sigmoid =
+                (z_val >= 0.0f) ? 1.0f / (1.0f + std::exp(-z_val))
+                                : std::exp(z_val) / (1.0f + std::exp(z_val));
+            y_val *= z_val * sigmoid;
+          }
+
+          out_h[d] = static_cast<input_t>(y_val);
+        }
+      }
+    }
+  }
+}
+
+}  // namespace mamba_cpu
@@ -213,6 +213,32 @@ void compute_slot_mapping_kernel_impl(const torch::Tensor query_start_loc,
                                      torch::Tensor slot_mapping,
                                      const int64_t block_size);

+at::Tensor causal_conv1d_update_cpu_impl(
+    at::Tensor& x, at::Tensor& conv_state, const at::Tensor& weight,
+    const c10::optional<at::Tensor>& bias,
+    const c10::optional<std::string>& activation,
+    const c10::optional<at::Tensor>& conv_state_indices,
+    const c10::optional<at::Tensor>& query_start_loc, int64_t pad_slot_id);
+
+void selective_state_update_cpu_impl(
+    at::Tensor& state, const at::Tensor& x, const at::Tensor& dt,
+    const at::Tensor& A, const at::Tensor& B, const at::Tensor& C,
+    const c10::optional<at::Tensor>& D, const c10::optional<at::Tensor>& z,
+    const c10::optional<at::Tensor>& dt_bias, bool dt_softplus,
+    const c10::optional<at::Tensor>& state_batch_indices,
+    const c10::optional<at::Tensor>& dst_state_batch_indices,
+    int64_t null_block_id, at::Tensor& out,
+    const c10::optional<at::Tensor>& num_accepted_tokens,
+    const c10::optional<at::Tensor>& cu_seqlens);
+
+void mamba_chunk_scan_fwd_cpu_impl(at::Tensor& out, at::Tensor& final_states,
+                                   const at::Tensor& x, const at::Tensor& dt,
+                                   const at::Tensor& A, const at::Tensor& B,
+                                   const at::Tensor& C,
+                                   const c10::optional<at::Tensor>& D,
+                                   const c10::optional<at::Tensor>& z,
+                                   const at::Tensor& cu_seqlens);
+
 void init_cpu_memory_env(std::vector<int64_t> node_ids);

 namespace cpu_utils {
@@ -595,6 +621,30 @@ TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops) {
      "block_size) -> ()",
      &compute_slot_mapping_kernel_impl);

+  // Mamba CPU kernels
+  ops.def(
+      "causal_conv1d_update_cpu_vec("
+      "Tensor(a0!) x, Tensor(a1!) conv_state, Tensor weight, "
+      "Tensor? bias, str? activation, Tensor? conv_state_indices, "
+      "Tensor? query_start_loc, SymInt pad_slot_id) -> Tensor",
+      &causal_conv1d_update_cpu_impl);
+
+  ops.def(
+      "selective_state_update_cpu("
+      "Tensor(a0!) state, Tensor x, Tensor dt, Tensor A, Tensor B, Tensor C, "
+      "Tensor? D, Tensor? z, Tensor? dt_bias, bool dt_softplus, "
+      "Tensor? state_batch_indices, Tensor? dst_state_batch_indices, "
+      "SymInt null_block_id, Tensor(a13!) out, "
+      "Tensor? num_accepted_tokens, Tensor? cu_seqlens) -> ()",
+      &selective_state_update_cpu_impl);
+
+  ops.def(
+      "mamba_chunk_scan_fwd_cpu("
+      "Tensor(a0!) out, Tensor(a1!) final_states, "
+      "Tensor x, Tensor dt, Tensor A, Tensor B, Tensor C, "
+      "Tensor? D, Tensor? z, Tensor cu_seqlens) -> ()",
+      &mamba_chunk_scan_fwd_cpu_impl);
+
  ops.def("init_cpu_memory_env(SymInt[] node_ids) -> ()", &init_cpu_memory_env);

  // Speculative decoding kernels
@@ -1174,7 +1174,8 @@ __global__ void cp_gather_and_upconvert_fp8_kv_cache(
    const int32_t num_reqs, const int32_t block_size,
    const int32_t total_tokens, 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 int64_t dst_entry_stride,
+    const int32_t* __restrict__ seq_starts) {  // Optional source offsets
  const int flat_warp_id = (blockIdx.x * blockDim.x + threadIdx.x) >> 5;
  if (flat_warp_id >= total_tokens) return;
  const int lane_id = threadIdx.x & 31;
@@ -1192,7 +1193,8 @@ __global__ void cp_gather_and_upconvert_fp8_kv_cache(

  // Compute physical token address via block table
  const int out_token_id = flat_warp_id;
-  const int token_offset = out_token_id - workspace_starts[req_id];
+  int token_offset = out_token_id - workspace_starts[req_id];
+  if (seq_starts != nullptr) token_offset += seq_starts[req_id];
  const int cache_block_idx = token_offset / block_size;
  const int offset_in_block = token_offset % block_size;
  const int physical_block =
@@ -1383,9 +1385,9 @@ void cp_gather_and_upconvert_fp8_kv_cache(
    torch::stable::Tensor const& src_cache,    // [NUM_BLOCKS, BLOCK_SIZE, 656]
    torch::stable::Tensor const& dst,          // [TOT_TOKENS, 576]
    torch::stable::Tensor const& block_table,  // [BATCH, BLOCK_INDICES]
-    torch::stable::Tensor const& seq_lens,     // [BATCH]
    torch::stable::Tensor const& workspace_starts,  // [BATCH]
-    int64_t batch_size) {
+    int64_t batch_size,
+    std::optional<torch::stable::Tensor> seq_starts = std::nullopt) {
  torch::stable::accelerator::DeviceGuard device_guard(
      src_cache.get_device_index());
  const cudaStream_t stream = get_current_cuda_stream();
@@ -1396,20 +1398,25 @@ void cp_gather_and_upconvert_fp8_kv_cache(
  STD_TORCH_CHECK(
      block_table.scalar_type() == torch::headeronly::ScalarType::Int,
      "block_table must be int32");
-  STD_TORCH_CHECK(seq_lens.scalar_type() == torch::headeronly::ScalarType::Int,
-                  "seq_lens must be int32");
  STD_TORCH_CHECK(
      workspace_starts.scalar_type() == torch::headeronly::ScalarType::Int,
      "workspace_starts must be int32");
+  if (seq_starts.has_value()) {
+    STD_TORCH_CHECK(
+        seq_starts.value().scalar_type() == torch::headeronly::ScalarType::Int,
+        "seq_starts must be int32");
+  }

  STD_TORCH_CHECK(src_cache.device() == dst.device(),
                  "src_cache and dst must be on the same device");
  STD_TORCH_CHECK(src_cache.device() == block_table.device(),
                  "src_cache and block_table must be on the same device");
-  STD_TORCH_CHECK(src_cache.device() == seq_lens.device(),
-                  "src_cache and seq_lens must be on the same device");
  STD_TORCH_CHECK(src_cache.device() == workspace_starts.device(),
                  "src_cache and workspace_starts must be on the same device");
+  if (seq_starts.has_value()) {
+    STD_TORCH_CHECK(src_cache.device() == seq_starts.value().device(),
+                    "src_cache and seq_starts must be on the same device");
+  }
  auto dtype = src_cache.scalar_type();
  STD_TORCH_CHECK(
      dtype == torch::headeronly::ScalarType::Byte ||               // uint8
@@ -1438,6 +1445,9 @@ void cp_gather_and_upconvert_fp8_kv_cache(
  constexpr int warps_per_block = 8;
  const int grid_size = (total_tokens + warps_per_block - 1) / warps_per_block;
  const int block_size_threads = warps_per_block * 32;  // 256 threads
+  const int32_t* seq_starts_ptr =
+      seq_starts.has_value() ? seq_starts.value().const_data_ptr<int32_t>()
+                             : nullptr;

  vllm::cp_gather_and_upconvert_fp8_kv_cache<<<grid_size, block_size_threads, 0,
                                               stream>>>(
@@ -1446,7 +1456,7 @@ void cp_gather_and_upconvert_fp8_kv_cache(
      workspace_starts.const_data_ptr<int32_t>(),
      static_cast<int32_t>(batch_size), block_size, total_tokens,
      block_table_stride, cache_block_stride, cache_entry_stride,
-      dst_entry_stride);
+      dst_entry_stride, seq_starts_ptr);
 }

 // Macro to dispatch the kernel based on the data type.
@@ -71,6 +71,73 @@ __device__ __forceinline__ float toFloat(T value) {
  }
 }

+#ifndef USE_ROCM
+// Adapted from:
+// https://github.com/sgl-project/sglang/blob/main/python/sglang/jit_kernel/csrc/deepseek_v4/hash_topk.cuh
+template <typename OutIndType, typename HashIndType>
+__launch_bounds__(128) __global__
+    void dsv4HashTopkSoftplusSqrt(const float* input, float* output,
+                                  OutIndType* indices, int num_rows,
+                                  int num_experts, float routed_scaling_factor,
+                                  const HashIndType* input_ids,
+                                  const HashIndType* tid2eid) {
+  const int warp = (blockIdx.x * blockDim.x + threadIdx.x) / 32;
+  const int lane = threadIdx.x % 32;
+  if (warp >= num_rows) return;
+  const int64_t token_id = load_index_as_int64(input_ids, warp);
+
+  #if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900)
+  cudaGridDependencySynchronize();
+  #endif
+  int expert = 0;
+  float weight = 0.f;
+  if (lane < 6) {
+    // only load and calculate for 6 experts
+    expert = static_cast<int>(tid2eid[token_id * 6 + lane]);
+    const float x = input[warp * num_experts + expert];
+    weight = sqrtf(fmaxf(x, 0.f) + __logf(1.f + __expf(-fabsf(x))));
+  }
+  float weight_sum = weight;
+  #pragma unroll
+  for (int mask = 16; mask > 0; mask >>= 1) {
+    // sum in warp
+    weight_sum += VLLM_SHFL_XOR_SYNC(weight_sum, mask);
+  }
+
+  #if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900)
+  cudaTriggerProgrammaticLaunchCompletion();
+  #endif
+  if (lane < 6) {
+    const int offset = warp * 6 + lane;
+    output[offset] =
+        weight * routed_scaling_factor / (weight_sum > 0.f ? weight_sum : 1.f);
+    indices[offset] = static_cast<OutIndType>(expert);
+  }
+}
+
+template <typename OutIndType, typename HashIndType>
+void launchDsv4HashTopk(const float* input, float* output, OutIndType* indices,
+                        int num_rows, int num_experts,
+                        double routed_scaling_factor,
+                        const HashIndType* input_ids,
+                        const HashIndType* tid2eid, cudaStream_t stream) {
+  if (num_rows == 0) return;
+  auto* kernel = &dsv4HashTopkSoftplusSqrt<OutIndType, HashIndType>;
+  cudaLaunchConfig_t config = {};
+  config.gridDim = (num_rows + 3) / 4;
+  config.blockDim = 128;
+  config.stream = stream;
+  cudaLaunchAttribute attr;
+  attr.id = cudaLaunchAttributeProgrammaticStreamSerialization;
+  attr.val.programmaticStreamSerializationAllowed = 1;
+  config.attrs = &attr;
+  config.numAttrs = 1;
+  const float scale = static_cast<float>(routed_scaling_factor);
+  cudaLaunchKernelEx(&config, kernel, input, output, indices, num_rows,
+                     num_experts, scale, input_ids, tid2eid);
+}
+#endif
+
 // ====================== TopK softplus_sqrt things
 // ===============================

@@ -556,6 +623,17 @@ void topkGatingSoftplusSqrtKernelLauncher(
    const float* correction_bias, const bool use_hash,
    const HashIndType* input_ids, const HashIndType* tid2eid,
    cudaStream_t stream) {
+#ifndef USE_ROCM
+  if constexpr (std::is_same_v<InputType, float>) {
+    if (use_hash && topk == 6 && renormalize &&
+        (num_experts == 256 || num_experts == 384)) {
+      launchDsv4HashTopk<IndType, HashIndType>(
+          gating_output, topk_weights, topk_indices, num_tokens, num_experts,
+          routed_scaling_factor, input_ids, tid2eid, stream);
+      return;
+    }
+  }
+#endif
  static constexpr int WARPS_PER_TB = 4;
  static constexpr int BYTES_PER_LDG_POWER_OF_2 = 16;
  // for bfloat16 dtype, we need 4 bytes loading to make sure num_experts
@@ -527,9 +527,9 @@ void cp_gather_and_upconvert_fp8_kv_cache(
                                                    // 656]
    torch::stable::Tensor const& dst,               // [TOT_TOKENS, 576]
    torch::stable::Tensor const& block_table,       // [BATCH, BLOCK_INDICES]
-    torch::stable::Tensor const& seq_lens,          // [BATCH]
    torch::stable::Tensor const& workspace_starts,  // [BATCH]
-    int64_t batch_size);
+    int64_t batch_size,
+    std::optional<torch::stable::Tensor> seq_starts = std::nullopt);

 // Indexer K quantization and cache function
 void indexer_k_quant_and_cache(
@@ -847,8 +847,8 @@ STABLE_TORCH_LIBRARY_FRAGMENT(_C_cache_ops, ops) {

  ops.def(
      "cp_gather_and_upconvert_fp8_kv_cache(Tensor src_cache, Tensor! dst, "
-      "Tensor block_table, Tensor seq_lens, Tensor workspace_starts, int "
-      "batch_size) -> ()");
+      "Tensor block_table, Tensor workspace_starts, int batch_size, Tensor? "
+      "seq_starts) -> ()");

  ops.def(
      "indexer_k_quant_and_cache(Tensor k, Tensor! kv_cache, Tensor "
@@ -5,7 +5,7 @@ vLLM uses the following environment variables to configure the system:
 !!! warning
    Please note that `VLLM_PORT` and `VLLM_HOST_IP` set the port and ip for vLLM's **internal usage**. It is not the port and ip for the API server. If you use `--host $VLLM_HOST_IP` and `--port $VLLM_PORT` to start the API server, it will not work.

-    All environment variables used by vLLM are prefixed with `VLLM_`. **Special care should be taken for Kubernetes users**: please do not name the service as `vllm`, otherwise environment variables set by Kubernetes might conflict with vLLM's environment variables, because [Kubernetes sets environment variables for each service with the capitalized service name as the prefix](https://kubernetes.io/docs/concepts/services-networking/service/#environment-variables).
+    Most vLLM-specific environment variables are prefixed with `VLLM_` (a handful of standard names — for example `CUDA_VISIBLE_DEVICES`, `MAX_JOBS`, `S3_ACCESS_KEY_ID`/`S3_SECRET_ACCESS_KEY`/`S3_ENDPOINT_URL`, `DO_NOT_TRACK`, `NO_COLOR` — are also read directly when set). **Special care should be taken for Kubernetes users**: please do not name the service as `vllm`, otherwise environment variables set by Kubernetes might conflict with vLLM's environment variables, because [Kubernetes sets environment variables for each service with the capitalized service name as the prefix](https://kubernetes.io/docs/concepts/services-networking/service/#environment-variables).

 ```python
 --8<-- "vllm/envs.py:env-vars-definition"
@@ -6,7 +6,11 @@ vLLM maintains a per-commit wheel repository (commonly referred to as "nightly")

 ### Wheel Building

-Wheels are built in the `Release` pipeline (`.buildkite/release-pipeline.yaml`) after a PR is merged into the main branch, with multiple variants:
+Wheels are built in the `Release` pipeline
+(`.buildkite/release-pipeline.yaml`) after a PR is merged into the main branch.
+Regular builds produce the CUDA 13.0 wheels for x86_64 and aarch64. Additional
+wheel variants and ROCm builds can be unblocked on demand and run automatically
+when `NIGHTLY=1`:

 - **Backend variants**: `cpu` and `cuXXX` (e.g., `cu129`, `cu130`).
 - **Architecture variants**: `x86_64` and `aarch64`.
@@ -205,7 +205,7 @@ hardware and configuration.

 | Backend | Description | Dtypes | Compute Cap. | Notes |
 | ------- | ----------- | ------ | ------------ | ----- |
-| `FLASH_ATTN`‡ | FlashAttention varlen (FA2/FA3/FA4) | fp16, bf16 | Any | (qk_nope_head_dim=128, qk_rope_head_dim=64, v_head_dim=128) (FA2/FA3/FA4) or (qk_nope_head_dim=192, qk_rope_head_dim=64, v_head_dim=256) (FA2/FA3 only) |
+| `FLASH_ATTN`‡ | FlashAttention varlen (FA2/FA3/FA4) | fp16, bf16 | Any | (qk_nope_head_dim=128, qk_rope_head_dim=64, v_head_dim=128) (FA2/FA3/FA4) or (qk_nope_head_dim=64, qk_rope_head_dim=64, v_head_dim=128) (FA2/FA3/FA4) or (qk_nope_head_dim=192, qk_rope_head_dim=64, v_head_dim=256) (FA2/FA3 only) |
 | `TRTLLM_RAGGED` | TensorRT-LLM ragged attention | fp16, bf16 | 10.x | (qk_nope_head_dim=128, qk_rope_head_dim=64, v_head_dim=128) or (qk_nope_head_dim=192, qk_rope_head_dim=64, v_head_dim=256) only |
 | `FLASHINFER` | FlashInfer CUTLASS backend | fp16, bf16 | 10.x | (qk_nope_head_dim=128, qk_rope_head_dim=64, v_head_dim=128) only |
 | `TOKENSPEED_MLA` | | fp16, bf16 | 10.x | (qk_nope_head_dim=128, qk_rope_head_dim=64, v_head_dim=128) only |
@@ -306,7 +306,7 @@ Supported quantization scheme/hardware combinations:

 - Pass: [`vllm/compilation/passes/fusion/rms_quant_fusion.py`](https://github.com/vllm-project/vllm/blob/main/vllm/compilation/passes/fusion/rms_quant_fusion.py)
 - ROCm AITER pass: [`vllm/compilation/passes/fusion/rocm_aiter_fusion.py`](https://github.com/vllm-project/vllm/blob/main/vllm/compilation/passes/fusion/rocm_aiter_fusion.py)
- CUDA/HIP kernels: [`csrc/layernorm_quant_kernels.cu`](https://github.com/vllm-project/vllm/blob/main/csrc/layernorm_quant_kernels.cu)
+- CUDA/HIP kernels: [`csrc/libtorch_stable/layernorm_quant_kernels.cu`](https://github.com/vllm-project/vllm/blob/main/csrc/libtorch_stable/layernorm_quant_kernels.cu)

 ### SiLU+Mul + Quantization (`fuse_act_quant`)

@@ -332,7 +332,7 @@ Supported quantization scheme/hardware combinations:
 - Pass: [`vllm/compilation/passes/fusion/act_quant_fusion.py`](https://github.com/vllm-project/vllm/blob/main/vllm/compilation/passes/fusion/act_quant_fusion.py)
 - ROCm AITER pass: [`vllm/compilation/passes/fusion/rocm_aiter_fusion.py`](https://github.com/vllm-project/vllm/blob/main/vllm/compilation/passes/fusion/rocm_aiter_fusion.py)
 - CUDA/HIP kernels: [`csrc/quantization/`](https://github.com/vllm-project/vllm/blob/main/csrc/quantization/)
- Fused SiLU+Mul+BlockQuant kernel: [`csrc/quantization/fused_kernels/fused_silu_mul_block_quant.cu`](https://github.com/vllm-project/vllm/blob/main/csrc/quantization/fused_kernels/fused_silu_mul_block_quant.cu)
+- Fused SiLU+Mul+BlockQuant kernel: [`csrc/libtorch_stable/quantization/fused_kernels/fused_silu_mul_block_quant.cu`](https://github.com/vllm-project/vllm/blob/main/csrc/libtorch_stable/quantization/fused_kernels/fused_silu_mul_block_quant.cu)

 ### RMSNorm + Padding (`fuse_act_padding`)

@@ -107,6 +107,7 @@ Batch invariance has been tested and verified on the following models:
 - **Llama 3**: Llama3.1 and 3.2 series, `meta-llama/Llama-3.2-3B-Instruct` for example
 - **GPT-OSS**: `openai/gpt-oss-20b`, `openai/gpt-oss-120b`
 - **Mistral**: `mistralai/Mistral-7B-v0.3`
+- **Phi series**: `microsoft/Phi-3.5-mini-instruct`

 Other models may also work, but these have been explicitly validated. If you encounter issues with a specific model, please report them on the [GitHub issue tracker](https://github.com/vllm-project/vllm/issues/new/choose).

@@ -68,13 +68,14 @@ vllm serve <model> \
 | --- | --- | --- | --- | --- |
 | `spec_name` | no | `CPUOffloadingSpec` | both | Set to `TieringOffloadingSpec` for multi-tier. |
 | `cpu_bytes_to_use` | yes | — | both | Total bytes of host memory reserved for the CPU tier across all workers (not per-worker). |
-| `block_size` | no | GPU block size | both | Offloaded block size in tokens; must be a multiple of the GPU block size. |
+| `block_size` | no | GPU block size | both | Offloaded block size in tokens; must be a multiple of the GPU block size. Mutually exclusive with `blocks_per_chunk`. |
+| `blocks_per_chunk` | no | `1` | both | Offloaded chunk size in GPU blocks; must be > 0. Alternative to `block_size` for models whose KV cache groups have different block sizes. |
 | `eviction_policy` | no | `lru` | both | Primary tier policy: `lru` or `arc`. |
 | `store_threshold` | no | `0` | single-tier | Min lookups before a block is offloaded. Values ≥ 2 are rejected by `TieringOffloadingSpec`. |
 | `max_tracker_size` | no | `64000` | single-tier | Max entries in the lookup tracker. |
 | `secondary_tiers` | no | `[]` | multi-tier | List of secondary tier configs (see below). |
 | `offload_prompt_only` | no | `true` | both | If `true`, only prompt (prefill) blocks are offloaded; decode blocks are skipped. |
-| `self_describing_kv_events` | no | `false` | single-tier | Opt-in. When `true` *and* KV cache events are enabled (`--kv-events-config` with `enable_kv_cache_events`), the connector emits self-describing block-granular `BlockStored`/`BlockRemoved` payloads (constituent block hashes, whole-chunk `token_ids`, per-block `block_size`, parent hash, LoRA + group/cache-spec metadata) instead of the placeholder fallback, so external KV-event consumers can index offloaded blocks. Inert unless events are enabled. Currently rejected by `TieringOffloadingSpec`. Full-attention groups only; sliding-window/SSM groups keep the placeholder fallback. In chunk mode (`block_size` > GPU block size), overlapping chunks re-announce shared per-block hashes, so consumers must reference-count (deduplicate) repeated store/remove announcements. |
+| `self_describing_kv_events` | no | `false` | single-tier | Opt-in. When `true` *and* KV cache events are enabled (`--kv-events-config` with `enable_kv_cache_events`), the connector emits self-describing block-granular `BlockStored`/`BlockRemoved` payloads (constituent block hashes, whole-chunk `token_ids`, per-block `block_size`, parent hash, LoRA + group/cache-spec metadata) instead of the placeholder fallback, so external KV-event consumers can index offloaded blocks. Inert unless events are enabled. Currently rejected by `TieringOffloadingSpec`. Full-attention groups only; sliding-window/SSM groups keep the placeholder fallback. In chunk mode (`block_size` > GPU block size, or `blocks_per_chunk` > 1), overlapping chunks re-announce shared per-block hashes, so consumers must reference-count (deduplicate) repeated store/remove announcements. |
 | `spec_module_path` | no | — | both | Python import path for a custom `OffloadingSpec` not in the built-in registry. Required only when `spec_name` is not built-in (advanced). |

 ## Secondary Tiers
@@ -83,9 +84,11 @@ Each entry in `secondary_tiers` is a dict with a required `type` field plus tier

 The filesystem and object-store tiers can publish hash-only `BlockStored` KV events for blocks they successfully store, tagged with a stable per-tier `medium` (`FS` for the filesystem tier, `OBJ` for the object-store tier). Set `enable_kv_events: true` in the tier's entry to opt in; events are published only when KV cache events are also enabled globally via `--kv-events-config`.

+Set the optional `locality` tier field to `LOCAL` or `REMOTE` to describe the tier's storage location relative to the publishing vLLM instance. `LOCAL` marks storage local to that instance, while `REMOTE` marks storage that is not local to it. When the setting is omitted, locality is unspecified. vLLM does not infer it from the tier type, so an OBJ tier is not implicitly `REMOTE`. A KV event includes `locality` only when the tier explicitly configures it. This metadata describes the tier property without implying that a consumer can already route requests to its blocks.
+
 ### Filesystem (FS)

-The filesystem tier (`type: "fs"`) writes blocks to a directory on local storage.
+The filesystem tier (`type: "fs"`) writes blocks to a filesystem directory.

 | Key | Required | Default | Notes |
 | --- | --- | --- | --- |
@@ -94,6 +97,7 @@ The filesystem tier (`type: "fs"`) writes blocks to a directory on local storage
 | `n_read_threads` | no | `16` | Read-priority I/O threads (load path). |
 | `n_write_threads` | no | `16` | Write-priority I/O threads (store path). |
 | `enable_kv_events` | no | `false` | Publish `BlockStored` KV events (medium `FS`) for successfully stored blocks. Requires KV cache events to be enabled globally. |
+| `locality` | no | unspecified | `LOCAL` or `REMOTE` relative to the publishing vLLM instance. Included in the tier's KV events only when explicitly configured. |

 Each thread group prefers its own queue but pulls from the other when its primary queue is empty, so a write-heavy or read-heavy burst won't leave the off-priority queue waiting. Size the totals to your storage's effective concurrency.

@@ -134,6 +138,7 @@ The object-store tier (`type: "obj"`) offloads blocks to an S3-compatible object
 | `prefix` | no | `""` | Key prefix prepended to all object keys. |
 | `io_threads` | no | `4` | Number of NIXL OBJ backend I/O threads. |
 | `enable_kv_events` | no | `false` | Publish `BlockStored` KV events (medium `OBJ`) for successfully stored blocks. Requires KV cache events to be enabled globally. |
+| `locality` | no | unspecified | `LOCAL` or `REMOTE` relative to the publishing vLLM instance. Included in the tier's KV events only when explicitly configured; OBJ does not imply `REMOTE`. |

 `store_config` fields:

@@ -175,7 +180,7 @@ Rather than embedding `host`/`port` in each `secondary_tiers` entry, set them on

 - `cpu_bytes_to_use`: a bigger CPU tier means fewer trips to slower secondary tiers and a higher hit rate. The value is total across all workers, not per-worker. Leave headroom for the rest of the host workload.
 - For single-tier (CPU-only) setups, set `cpu_bytes_to_use` larger than the aggregate GPU KV cache. Because offloading is immediate, a smaller CPU tier just mirrors what the GPU already holds and adds no hit rate.
- `block_size`: larger offloaded blocks reduce per-block bookkeeping overhead but increase the granularity of lookups. Must be a multiple of the GPU block size.
+- `block_size` / `blocks_per_chunk`: larger offloaded chunks reduce per-block bookkeeping overhead but increase the granularity of lookups.
 - FS thread counts: tune `n_read_threads` and `n_write_threads` to the parallelism your storage can sustain. Reads are latency-sensitive on the prefill path, so prefer more read threads when prefill hit rates are high.
 - Sharing `root_dir` across runs: runs with the same model, `block_size`, parallelism layout, and dtype share files under the same `<digest>` subdirectory. Changing any of these produces a new subdirectory; old ones are orphaned but harmless. Delete them to reclaim disk.

@@ -1,5 +1,5 @@
 {% extends "base.html" %}

 {% block announce %}
-  <p>You are viewing the latest developer preview docs. <a href="https://docs.vllm.ai/en/stable/">Click here</a> to view docs for the latest stable release.</p>
+  <p>You are viewing the latest developer preview docs. <a href="https://docs.vllm.ai/en/stable/{{ page.url }}">Click here</a> to view docs for the latest stable release.</p>
 {% endblock %}
@@ -31,10 +31,8 @@
 | THUDM/CodeGeex4-All-9B                             | CodeGeexForCausalLM                              | ✅                    |                        |       |
 | chuhac/TeleChat2-35B                               | LlamaForCausalLM (TeleChat2 based on Llama arch) | ✅                    |                        |       |
 | 01-ai/Yi1.5-34B-Chat                               | YiForCausalLM                                    | ✅                    |                        |       |
-| THUDM/CodeGeex4-All-9B                             | CodeGeexForCausalLM                              | ✅                    |                        |       |
 | deepseek-ai/DeepSeek-Coder-33B-base                | DeepSeekCoderForCausalLM                         | ✅                    |                        |       |
 | meta-llama/Llama-2-13b-chat-hf                     | LlamaForCausalLM                                 | ✅                    |                        |       |
-| THUDM/CodeGeex4-All-9B                             | CodeGeexForCausalLM                              | ✅                    |                        |       |
 | Qwen/Qwen1.5-14B-Chat                              | QwenForCausalLM                                  | ✅                    |                        |       |
 | Qwen/Qwen1.5-32B-Chat                              | QwenForCausalLM                                  | ✅                    |                        |       |
 | RedHatAI/Meta-Llama-3.1-8B-Instruct-FP8-dynamic    | LlamaForCausalLM                                 |                       | ✅                     |       |
@@ -220,8 +220,6 @@ For multi-node deployment, add these EPLB flags to each node's command. We recom

 - Use simulator flags `VLLM_MOE_ROUTING_SIMULATION_STRATEGY=uniform_random` and `VLLM_RANDOMIZE_DP_DUMMY_INPUTS=1` so token routing is balanced across EP ranks.

- Increasing `VLLM_MOE_DP_CHUNK_SIZE` may increase throughput by increasing the maximum batch size for inter-rank token transfers. This may cause DeepEP  to throw `assert self.nvshmem_qp_depth >= (num_max_dispatch_tokens_per_rank + 1) * 2`, which can be fixed by increasing environment variable `NVSHMEM_QP_DEPTH`.
-
 ## Disaggregated Serving (Prefill/Decode Split)

 For production deployments requiring strict SLA guarantees for time-to-first-token and inter-token latency, disaggregated serving allows independent scaling of prefill and decode operations.
@@ -137,8 +137,12 @@ For further details on renderer APIs, please refer to [this page](renderer.md).

 ### Derenderer APIs

- `/v1/completions/derender` - Derenderer completion requests
- `/v1/chat/completions/derender` - Derenderer chat completion requests
+For further details on derenderer APIs, please refer to [this page](derenderer.md).
+
+- [Chat Completions Derender API](derenderer.md) (`/v1/chat/completions/derender`)
+    - Derender chat completion requests
+- [Completions Derender API](derenderer.md) (`/v1/completions/derender`)
+    - Derender completion requests

 ## Tokenize APIs

@@ -0,0 +1,98 @@
+# Derenderer APIs
+
+The derenderer API is the post processing counterpart to the [Renderer APIs](renderer.md). Where `/render` turns a request into token ID (preprocessing), `/derender` turns generated token IDs back into a fully formed OpenAI compatible response (detokenization, reasoning parsing, tool call parsing), all without a GPU.
+
+This closes the loop for a token-in / token-out engine in disaggregated serving:
+
+- **GPU less post processing**: Detokenization, reasoning parsing, and tool call parsing run on the same GPU less frontend that hosts `/render`
+- **Parser parity**: The derenderer reuses vLLM's tool and reasoning parsers, so a disaggregated deployment produces the same `content`/`reasoning`/ `tool_calls` split as a standard `vllm serve` server
+- **Non-streaming**: The endpoints expect a complete `GenerateResponse` with all token IDs present and perform one-shot parsing. Streaming derender would require a separate endpoint design and is not currently supported but is in the pipeline
+
+Both endpoints are hosted by the GPU less rendering server started with [`vllm launch render`](../../cli/launch/render.md), alongside the `/render`
+endpoints.
+
+## Pipeline
+
+```text
+                render                 generate                derender
+  request  ───────────────▶  token_ids  ─────────▶  token_ids  ──────────▶  response
+ (chat /            (GPU less)          (token-in /            (GPU less)   (OpenAI
+ completion)            │               token-out engine)          ▲        compatible)
+                        └─────────────── request + prompt_tokens ──┘
+```
+
+The derender step needs more than the engine's `token_ids`. It also consumes the original `chat_request`/`completion_request` and `prompt_tokens` carried over from the render step (see [Request format](#request-format)) so the tool and reasoning parsers have the context they need.
+
+## API Reference
+
+- Chat Completions Derender API (`/v1/chat/completions/derender`)
+    - Post process a single `GenerateResponse` into a `ChatCompletionResponse`
+- Completions Derender API (`/v1/completions/derender`)
+    - Post process a list of `GenerateResponse` objects (one per prompt) into a `CompletionResponse`
+
+## Request format
+
+Each request wraps the engine's `GenerateResponse`(s) together with the caller metadata needed to reconstruct the final response without a GPU.
+
+`/v1/chat/completions/derender`:
+
+??? code
+
+    ```python
+    --8<-- "vllm/entrypoints/scale_out/token_in_token_out/protocol.py:derender-chat-request"
+    ```
+
+`/v1/completions/derender`:
+
+??? code
+
+    ```python
+    --8<-- "vllm/entrypoints/scale_out/token_in_token_out/protocol.py:derender-completion-request"
+    ```
+
+Oversized payloads are rejected with a `400` before any `tokenizer.decode()` or parser runs.
+
+## Example
+
+The example below drives the full `render → generate → derender` round trip for a chat request against a GPU less render server (`/render`, `/derender`) and a token-in / token-out engine (`/inference/v1/generate`).
+
+```python
+import httpx
+
+MODEL = "meta-llama/Llama-3.2-1B-Instruct"
+RENDER = "http://localhost:8100"  # vllm launch render ...
+ENGINE = "http://localhost:8200"  # token-in / token-out engine
+
+chat_request = {
+    "model": MODEL,
+    "messages": [{"role": "user", "content": "What is 2+2?"}],
+    "max_tokens": 32,
+}
+
+with httpx.Client(timeout=60.0) as client:
+    # 1. Render: request -> token IDs (GPU less)
+    generate_request = client.post(
+        f"{RENDER}/v1/chat/completions/render", json=chat_request
+    ).json()
+    prompt_tokens = len(generate_request["token_ids"])
+
+    # 2. Generate: token IDs -> token IDs (token-in / token-out engine)
+    generate_response = client.post(
+        f"{ENGINE}/inference/v1/generate", json=generate_request
+    ).json()
+
+    # 3. Derender: token IDs -> ChatCompletionResponse (GPU less)
+    response = client.post(
+        f"{RENDER}/v1/chat/completions/derender",
+        json={
+            "model": MODEL,
+            "generate_response": generate_response,
+            "prompt_tokens": prompt_tokens,
+            "chat_request": chat_request,
+        },
+    ).json()
+
+print(response["choices"][0]["message"]["content"])
+```
+
+Passing `chat_request` lets the derenderer run the configured tool and reasoning parsers. This means `response["choices"][0]["message"]` carries the same `content` / `reasoning` / `tool_calls` split a `vllm serve` server would produce. Omit `chat_request` for plain detokenization only.
@@ -12,3 +12,5 @@ Our renderer API is designed to disaggregate the render phase(preprocessing) and
    - Render completion requests
 - [Chat Completions Render API](renderer.md) (`/v1/chat/completions/render`)
    - Render chat completions
+
+For the post processing counterpart that turns generated token IDs back into OpenAI compatible responses, see the [Derenderer APIs](derenderer.md).
@@ -155,8 +155,10 @@ When `--api-key` is configured, the following `/v1` endpoints require Bearer tok
 - `/v1/chat/completions` - Chat completions
 - `/v1/chat/completions/batch` - Batch chat completions
 - `/v1/chat/completions/render` - Render chat completion requests
+- `/v1/chat/completions/derender` - Derender chat completion requests
 - `/v1/completions` - Text completions
 - `/v1/completions/render` - Render completion requests
+- `/v1/completions/derender` - Derender completion requests
 - `/v1/embeddings` - Generate embeddings
 - `/v1/audio/transcriptions` - Audio transcription
 - `/v1/audio/translations` - Audio translation
@@ -210,8 +210,31 @@ async def stream_decode_response(session, response, request_id):
        await session.close()


-def example_round_robin_dp_loader(request_number, dp_size):
-    return request_nums % dp_size
+def flat_interleaved_dp_route(request_number, instances):
+    """Flat round-robin over the full (instance, dp_rank) slot space.
+
+    ONE counter over (n_instances * dp_size) slots, so instance-selection and
+    DP-rank-selection are derived from the SAME index and can never alias. The
+    previous scheme computed instance = req % n and rank = req % dp from the
+    same counter with n | dp, which locked each instance to a stride-n subset
+    of its ranks (e.g. 2 prefill instances -> 4 of 8 ranks each -> half the
+    GPUs never receive a request, so the deployment falsely appears not to
+    scale).
+
+    Interleaved order — inst0_r0, inst1_r0, inst0_r1, inst1_r1, ... — so
+    consecutive requests alternate instances AND every rank gets walked.
+
+    Assumes homogeneous dp_size across a role's instances (true for the
+    DP<->DP and DP<->TP deployments this proxy targets). Returns
+    (instance_index, dp_rank); dp_rank is None when dp_size == 1 (e.g. a TP
+    decode), which avoids forwarding an out-of-range data-parallel rank.
+    """
+    n = len(instances)
+    dp = instances[0]["dp_size"]
+    slot = (request_number - 1) % (n * dp)
+    inst_idx = slot % n
+    dp_rank = (slot // n) if dp > 1 else None
+    return inst_idx, dp_rank


@app.route("/health", methods=["GET"])
@@ -252,18 +275,21 @@ async def handle_request(api: str, request: Request):
                    503,
                )
            )
-        pid = request_nums % len(prefill_instances)
-        did = request_nums % len(decode_instances)
+        # Flat interleaved round-robin (see flat_interleaved_dp_route): ONE
+        # counter over the full (instance, dp_rank) slot space per role, so
+        # instance-selection and DP-rank-selection derive from the same index
+        # and can never alias. The old scheme keyed both on request_nums with
+        # n_instances | dp_size, stranding half the ranks (e.g. in 2P_DP8EP).
+        pid, selected_prefill_dp_rank = flat_interleaved_dp_route(
+            request_nums, prefill_instances
+        )
+        # Decode instance selection uses the same interleaved walk; in READ
+        # mode the decode reads KV from selected_prefill_dp_rank, so the
+        # decode's own dp_rank is not forwarded here.
+        did, _ = flat_interleaved_dp_route(request_nums, decode_instances)
        prefill_instance_endpoint = prefill_instances[pid]
        decode_instance_endpoint = decode_instances[did]

-        selected_prefill_dp_rank = None
-        if prefill_instance_endpoint["dp_size"] > 1:
-            selected_prefill_dp_rank = example_round_robin_dp_loader(
-                request_nums // len(prefill_instance_endpoint),
-                prefill_instance_endpoint["dp_size"],
-            )
-
        # Embed both zmq_addresses in the request_id so the connector can parse
        # the peer's host/ports from it, similar to P2P-NCCL
        uid = str(uuid.uuid4()).replace("-", "")
@@ -427,9 +453,33 @@ if __name__ == "__main__":
    args = parser.parse_args()

    t = start_service_discovery("0.0.0.0", 36367)
-    app.debug = True
+    # High-concurrency hardening. Quart's app.run() uses a shallow listen
+    # backlog (100) and, with app.debug=True, adds per-request overhead that
+    # starves the single accept loop. Under a burst of ~512 simultaneous client
+    # connections the backlog overflows and the kernel RSTs the excess, so
+    # clients see "ClientOSError: [Errno 104] Connection reset by peer" before
+    # any response (~16% request loss at c=512). Serve via hypercorn with debug
+    # OFF and a deep backlog so the burst QUEUES (higher TTFT) instead of being
+    # reset -> 100% request success.
+    app.debug = False
    app.config["BODY_TIMEOUT"] = 360000
    app.config["RESPONSE_TIMEOUT"] = 360000

-    app.run(host="0.0.0.0", port=args.port)
+    import asyncio
+    import os
+
+    from hypercorn.asyncio import serve as _hypercorn_serve
+    from hypercorn.config import Config as _HypercornConfig
+
+    _hcfg = _HypercornConfig()
+    _hcfg.bind = [f"0.0.0.0:{args.port}"]
+    # Deep listen backlog so a wide connection burst queues, not RSTs. NOTE:
+    # effective backlog is capped by the host's net.core.somaxconn (proxy runs
+    # --network host); kernel 6.x defaults to 4096. Override via
+    # PROXY_LISTEN_BACKLOG.
+    _hcfg.backlog = int(os.environ.get("PROXY_LISTEN_BACKLOG", "4096"))
+    # Long-lived SSE streams (8k1k decode ~5 min): never reap on keepalive.
+    _hcfg.keep_alive_timeout = 360000.0
+
+    asyncio.run(_hypercorn_serve(app, _hcfg))
    t.join()
@@ -42,12 +42,16 @@ class BlockStored(KVCacheEvent):
    """

    group_idx: int | None = None
+    kv_cache_spec_kind: str | None = None
+    kv_cache_spec_sliding_window: int | None = None
+    locality: str | None = None


 class BlockRemoved(KVCacheEvent):
    block_hashes: list[ExternalBlockHash]
    medium: str | None
    group_idx: int | None = None
+    locality: str | None = None


 class AllBlocksCleared(KVCacheEvent):
@@ -17,6 +17,7 @@ from transformers import AutoProcessor, AutoTokenizer
 from vllm import LLM, EngineArgs, SamplingParams
 from vllm.lora.request import LoRARequest
 from vllm.multimodal.utils import fetch_image
+from vllm.platforms import current_platform
 from vllm.utils.argparse_utils import FlexibleArgumentParser

 QUESTION = "What is the content of each image?"
@@ -1443,6 +1444,8 @@ def run_generate(
    engine_args.seed = seed
    if tensor_parallel_size is not None:
        engine_args.tensor_parallel_size = tensor_parallel_size
+    if current_platform.is_rocm():
+        os.environ["VLLM_WORKER_MULTIPROC_METHOD"] = "spawn"
    llm = LLM.from_engine_args(engine_args)

    sampling_params = SamplingParams(
@@ -1484,6 +1487,8 @@ def run_chat(
    engine_args.seed = seed
    if tensor_parallel_size is not None:
        engine_args.tensor_parallel_size = tensor_parallel_size
+    if current_platform.is_rocm():
+        os.environ["VLLM_WORKER_MULTIPROC_METHOD"] = "spawn"
    llm = LLM.from_engine_args(engine_args)

    sampling_params = (
@@ -21,6 +21,7 @@ from vllm.assets.image import ImageAsset
 from vllm.assets.video import VideoAsset
 from vllm.lora.request import LoRARequest
 from vllm.multimodal.image import convert_image_mode
+from vllm.platforms import current_platform
 from vllm.utils.argparse_utils import FlexibleArgumentParser


@@ -2646,6 +2647,8 @@ def main(args):
    if args.tensor_parallel_size is not None:
        engine_args.tensor_parallel_size = args.tensor_parallel_size
    engine_args = maybe_add_vit_cuda_graph_compilation_config(args, engine_args)
+    if current_platform.is_rocm():
+        os.environ["VLLM_WORKER_MULTIPROC_METHOD"] = "spawn"
    llm = LLM.from_engine_args(engine_args)

    # Don't want to check the flag multiple times, so just hijack `prompts`.
@@ -122,7 +122,8 @@ python = "./.venv"
 [tool.typos.files]
 # these files may be written in non english words
 extend-exclude = ["tests/models/fixtures/*", "tests/prompts/*", "tests/tokenizers_/*",
-    "benchmarks/sonnet.txt", "tests/lora/data/*", "build/*",
+    "benchmarks/sonnet.txt", "rust/src/bench/src/datasets/sonnet.txt",
+    "tests/lora/data/*", "build/*",
    "examples/pooling/token_embed/*", "tests/models/language/pooling/*",
    "vllm/third_party/*", "vllm/entrypoints/serve/instrumentator/static/*",
    "tests/entrypoints/speech_to_text/transcription/test_transcription_validation.py",
@@ -379,7 +379,7 @@ inflect==5.6.2
    # via datamodel-code-generator
 iniconfig==2.0.0
    # via pytest
-instanttensor==0.1.5
+instanttensor==0.1.9
    # via -r requirements/test/cuda.in
 interegular==0.3.3
    # via lm-format-enforcer
@@ -58,7 +58,7 @@ arctic-inference == 0.1.1; platform_machine == "x86_64" # Required for suffix de
 numba == 0.65.0 # Required for N-gram speculative decoding
 runai-model-streamer[s3,gcs,azure]==0.15.7
 fastsafetensors>=0.3.2
-instanttensor>=0.1.5; platform_machine == "x86_64"
+instanttensor>=0.1.9; platform_machine == "x86_64"
 decord==0.6.0; platform_machine == "x86_64"
 # terratorch is temporarily disabled while PyPI has the `lightning` package
 # in `quarantined` status (every published terratorch version transitively
@@ -398,7 +398,7 @@ inflect==5.6.2
    # via datamodel-code-generator
 iniconfig==2.0.0
    # via pytest
-instanttensor==0.1.5
+instanttensor==0.1.9
    # via -r requirements/test/cuda.in
 interegular==0.3.3
    # via lm-format-enforcer
@@ -44,5 +44,5 @@ numba == 0.65.0 # Required for N-gram speculative decoding
 numpy
 runai-model-streamer[s3,gcs,azure]==0.15.7
 fastsafetensors>=0.3.2
-instanttensor>=0.1.5
+instanttensor>=0.1.9
 pydantic>=2.12 # 2.11 leads to error on python 3.13
@@ -54,7 +54,7 @@ arctic-inference==0.1.1 # Required for suffix decoding test
 numba==0.65.0 # Required for N-gram speculative decoding
 runai-model-streamer[s3,gcs,azure]==0.15.7
 fastsafetensors>=0.3.2
-instanttensor>=0.1.5
+instanttensor>=0.1.9
 decord==0.6.0

 # Prithvi tests
@@ -391,7 +391,7 @@ inflect==7.5.0
    # via datamodel-code-generator
 iniconfig==2.3.0
    # via pytest
-instanttensor==0.1.6
+instanttensor==0.1.9
    # via -r requirements/test/rocm.in
 interegular==0.3.3
    # via lm-format-enforcer
@@ -18,4 +18,4 @@ torchvision
 torchcodec >= 0.14 # Required for the torchcodec video decoding backend

 auto_round_lib==0.14.1
-vllm_xpu_kernels @ https://github.com/vllm-project/vllm-xpu-kernels/releases/download/v0.1.11/vllm_xpu_kernels-0.1.11-cp38-abi3-manylinux_2_28_x86_64.whl
+vllm_xpu_kernels @ https://github.com/vllm-project/vllm-xpu-kernels/releases/download/v0.1.11.1/vllm_xpu_kernels-0.1.11.1-cp38-abi3-manylinux_2_28_x86_64.whl
@@ -2585,6 +2585,7 @@ source = "registry+https://github.com/rust-lang/crates.io-index"
 checksum = "071dfc062690e90b734c0b2273ce72ad0ffa95f0c74596bc250dcfd960262841"
 dependencies = [
 "autocfg",
+ "libm",
 ]

 [[package]]
@@ -3305,6 +3306,16 @@ dependencies = [
 "getrandom 0.3.4",
 ]

+[[package]]
+name = "rand_distr"
+version = "0.5.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "6a8615d50dcf34fa31f7ab52692afec947c4dd0ab803cc87cb3b0b4570ff7463"
+dependencies = [
+ "num-traits",
+ "rand 0.9.2",
+]
+
 [[package]]
 name = "rawpointer"
 version = "0.2.1"
@@ -3435,6 +3446,7 @@ dependencies = [
 "base64 0.22.1",
 "bytes",
 "encoding_rs",
+ "futures-channel",
 "futures-core",
 "futures-util",
 "h2",
@@ -3549,6 +3561,15 @@ dependencies = [
 "rustc-hash 2.1.1",
 ]

+[[package]]
+name = "rlimit"
+version = "0.11.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "f35ee2729c56bb610f6dba436bf78135f728b7373bdffae2ec815b2d3eb98cc3"
+dependencies = [
+ "libc",
+]
+
 [[package]]
 name = "rmp"
 version = "0.8.15"
@@ -5439,6 +5460,38 @@ version = "0.9.5"
 source = "registry+https://github.com/rust-lang/crates.io-index"
 checksum = "0b928f33d975fc6ad9f86c8f283853ad26bdd5b10b7f1542aa2fa15e2289105a"

+[[package]]
+name = "vllm-bench"
+version = "0.1.0"
+dependencies = [
+ "anyhow",
+ "base64 0.22.1",
+ "bytes",
+ "chrono",
+ "clap",
+ "dirs",
+ "futures",
+ "hf-hub",
+ "image",
+ "indicatif",
+ "mimalloc",
+ "rand 0.9.2",
+ "rand_distr",
+ "rayon",
+ "reqwest 0.12.28",
+ "rlimit",
+ "rustc-hash 1.1.0",
+ "serde",
+ "serde_json",
+ "thiserror 2.0.18",
+ "tiktoken-rs 0.9.1",
+ "tokenizers",
+ "tokio",
+ "tokio-stream",
+ "url",
+ "uuid",
+]
+
 [[package]]
 name = "vllm-chat"
 version = "0.1.0"
@@ -5507,6 +5560,7 @@ dependencies = [
 "tracing",
 "tracing-subscriber",
 "uuid",
+ "vllm-bench",
 "vllm-chat",
 "vllm-engine-core-client",
 "vllm-managed-engine",
@@ -1,5 +1,6 @@
 [workspace]
 members = [
+    "src/bench",
    "src/chat",
    "src/cmd",
    "src/engine-core-client",
@@ -32,8 +33,10 @@ base64 = "0.22.1"
 bytemuck = { version = "1.25.0", features = ["extern_crate_alloc"] }
 byteorder = "1.5.0"
 bytes = "1.12.0"
+chrono = "0.4.42"
 clap = { version = "4.5.38", features = ["derive", "env"] }
 criterion = "0.5.1"
+dirs = "6.0.0"
 easy-ext = "1.0.3"
 educe = "0.6.0"
 enum-as-inner = "0.7.0"
@@ -50,7 +53,9 @@ hyper-util = { version = "0.1.20", features = [
    "service",
    "tokio",
 ] }
+image = { version = "0.25.9", default-features = false, features = ["jpeg"] }
 indexmap = "2.13.0"
+indicatif = "0.18.4"
 itertools = "0.14.0"
 libc = "0.2.177"
 llm-multimodal = { git = "https://github.com/smg-project/llm-multimodal", rev = "5390032d6dc8a3e6fdc83acd320260367eb4b9b5", default-features = false, features = ["native-tls"] }
@@ -71,10 +76,13 @@ prost-types = "0.14.3"
 pyo3 = "0.28.3"
 pythonize = "0.28.0"
 rand = "0.9.2"
+rand_distr = "0.5.1"
+rayon = "1.11.0"
 reasoning-parser = "1.2.2"
 reqwest = { version = "0.12.8", default-features = false, features = ["native-tls"] }
 reqwest-0-13 = { package = "reqwest", version = "0.13.4", default-features = false, features = ["native-tls"] }
 riptoken = { version = "0.3.0", default-features = false }
+rlimit = "0.11.0"
 rmp-serde = "1.3.1"
 rmpv = { version = "1.3.1", features = ["with-serde"] }
 rustc-hash = "1.1.0"
@@ -121,8 +129,10 @@ tracing = { version = "0.1.44", features = ["release_max_level_debug"] }
 tracing-futures = { version = "0.2.5", features = ["futures-03"] }
 tracing-subscriber = { version = "0.3.20", features = ["env-filter", "fmt"] }
 trait-set = "0.3.0"
+url = "2.5.7"
 uuid = { version = "1.22.0", features = ["v4"] }
 validator = { version = "0.20.0", features = ["derive"] }
+vllm-bench = { path = "src/bench" }
 vllm-chat = { path = "src/chat" }
 vllm-engine-core-client = { path = "src/engine-core-client" }
 vllm-llm = { path = "src/llm" }
@@ -0,0 +1,181 @@
+# AGENTS.md
+
+## Project Overview
+
+Rust rewrite of `vllm bench serve` — a high-performance benchmark client for vLLM serving endpoints. Standalone binary, no Python dependency at runtime.
+
+Member crate `vllm-bench` of the `rust/` workspace. Uses workspace dependencies and lints; the workspace `[profile.release]` (thin LTO, `panic = "abort"`) applies. Note the workspace bans rustls/ring (`rust/deny.toml`) — all HTTP must stay on native-tls, which is why HF Hub downloads go through `src/hub.rs` (async hf-hub API bridged to sync) instead of hf-hub's ureq backend.
+
+## Build & Test
+
+Run from the `rust/` workspace root:
+
+```bash
+# Build release binary (rust/target/release/vllm-bench)
+cargo build -p vllm-bench --release
+
+# Run all tests
+cargo test -p vllm-bench
+
+# Run ignored integration tests (requires network for tokenizer download)
+cargo test -p vllm-bench -- --ignored
+```
+
+## Architecture
+
+- `src/main.rs` — Entry point, mimalloc, tokio runtime, mode dispatch (compare/sweep/multi-run/multi-turn/single)
+- `src/cli.rs` — clap derive CLI args (~50+ flags)
+- `src/config.rs` — Validated config from CLI; `GoodputConfig`, `RampUpConfig`, sampling param merging
+- `src/error.rs` — `BenchError` enum (Http, Json, Tokenizer, Config, EndpointTimeout, Backend, Io)
+- `src/benchmark.rs` — Core benchmark orchestrator (spawn-per-request with tokio + Semaphore; fetches speculative decoding metrics from `/metrics`)
+- `src/multi_turn.rs` — Multi-turn conversation orchestrator (channel-based worker pool, sequential turns per conversation)
+- `src/sweep.rs` — Concurrency/rate parameter sweep (`--sweep-max-concurrency`, `--sweep-request-rate`)
+- `src/multi_run.rs` — N-run aggregation with mean/std/min/max/CV (`--num-runs`)
+- `src/compare.rs` — Side-by-side diff of two result JSON files (`--compare`)
+- `src/tokenizer.rs` — `TokenizerKind` enum: Local(HuggingFace), Tiktoken, OR Server-side `/tokenize`+`/detokenize` fallback
+- `src/tiktoken.rs` — Tiktoken BPE loader (`.tiktoken`/`.model` files; built-in encodings o200k_base/cl100k_base; pat_str extraction from Python source)
+- `src/hub.rs` — `HubRepo`: sync facade over hf-hub's async (reqwest/native-tls) API — per-download thread with its own runtime; the sync ureq backend is unusable here because it pulls rustls, which `rust/deny.toml` bans
+- `src/rate_control.rs` — Gamma/Poisson request scheduling + linear/exponential ramp-up
+- `src/ready_checker.rs` — Endpoint readiness with retry
+- `src/backends/` — Backend implementations (enum dispatch, not trait objects)
+    - `mod.rs` — `Backend` enum, `RequestFuncInput`/`RequestFuncOutput` (includes `messages` field for multi-turn)
+    - `streaming.rs` — SSE parser (`StreamedResponseHandler`) with speculative JSON parse for split TCP segments
+    - `openai_completions.rs` — `/v1/completions` backend
+    - `openai_chat.rs` — `/v1/chat/completions` backend (uses `input.messages` when set; zero-copy raw JSON payload for multimodal)
+    - `pooling.rs` — Non-streaming pooling/embedding backends: `openai-embeddings`, `openai-embeddings-chat`, `vllm-pooling`, `vllm-rerank`
+- `src/datasets/random.rs` — Random dataset generation with rayon parallelism
+- `src/datasets/random_mm.rs` — Random multimodal dataset (synthetic JPEG images, bucket config sampling, pre-serialized JSON fragments); `--enable-multimodal-chat` pre-builds the chat `messages` array at dataset time (mirrors Python's `apply_multimodal_chat_transformation`)
+- `src/datasets/sharegpt.rs` — ShareGPT JSON loader + HuggingFace Hub auto-download with caching
+- `src/datasets/sonnet.rs` and `src/datasets/sonnet.txt` — Sonnet dataset (built-in Shakespeare sonnets via `include_str!("sonnet.txt")`; controllable token length + shared prefix; mirrors Python `SonnetDataset`)
+- `src/datasets/speed_bench.rs` — NVIDIA SPEED-Bench loader (HF datasets-server API, 6 configs, 11 categories, local cache)
+- `src/datasets/hf_dataset.rs` — Generic HuggingFace dataset loader (datasets-server API, column auto-detection)
+- `src/datasets/custom.rs` — Custom JSONL dataset (`{"prompt": ..., "output_tokens": ...}` per line; `--custom-output-len -1` uses per-line output_tokens; prompts always sent raw — no client-side chat template)
+- `src/datasets/prefix_repetition.rs` — Prefix repetition dataset (N shared prefixes × fresh random suffixes, standard prefix-cache stress; mirrors Python `PrefixRepetitionRandomDataset`)
+- `src/datasets/random_rerank.rs` — Random rerank dataset (one query + batched documents per request for `vllm-rerank`; `--no-reranker` for embedding-based scoring; mirrors Python `RandomDatasetForReranking`)
+- `src/datasets/multi_turn.rs` — Multi-turn synthetic generator + ShareGPT multi-turn loader (3-tier prefix sharing: global/conversation/unique-suffix; `per_turn_input_len`)
+- `src/metrics/mod.rs` — `BenchmarkMetrics` and `MultiTurnMetrics` structs
+- `src/metrics/calculator.rs` — TTFT/TPOT/ITL/E2EL/throughput stats, goodput SLO checking, peak concurrency, `calculate_multi_turn_metrics`
+- `src/metrics/steady_state.rs` — Steady-state window detection (in-flight concurrency plateau via two-pointer start/end merge) + plateau throughput/TTFT/TPOT; gated on `--max-concurrency` set + `--request-rate inf` (closed-loop)
+- `src/output/console.rs` — Terminal output matching Python format + multi-turn per-turn breakdown
+- `src/output/json.rs` — JSON result file (compatible with Python schema) + multi-turn JSON with `per_turn_metrics`
+
+## Key Design Decisions
+
+- **Enum dispatch** for backends (avoids async trait object issues with `dyn`)
+- **reqwest http1_only()** to match Python aiohttp behavior
+- **rayon** for parallel dataset generation (key perf win over Python)
+- **mimalloc** global allocator to reduce contention at 1400+ concurrency (page-agnostic; works on aarch64 64K-page kernels where jemalloc aborts with `LG_PAGE=12` builds)
+- **Arc\<str\> prompts** zero-copy sharing across tokio tasks (~3GB savings at 100k prompts with 8k-token inputs)
+- **Spawn-per-request** `tokio::spawn` + `Semaphore` (matches Python asyncio pattern)
+- **Speculative JSON parse** in SSE handler — detects complete JSON before `\n\n` arrives, improving TTFT/ITL accuracy when TCP segments split
+- **Tokenizer fallback chain**: Local HF → Tiktoken (`.tiktoken`/`.model` + built-in encodings) → Server-side `/tokenize`+`/detokenize`. Blocking HTTP in rayon threads for server fallback.
+- **hf-hub** for downloading tokenizers and datasets from HuggingFace Hub
+- **Pre-serialized mm fragments** (`Arc<str>`) for multimodal: image content stored as JSON strings, zero-copy concatenated into payload — avoids deep-cloning ~200KB+ base64 per request
+- **Steady-state metrics** (default-on in closed-loop): measure throughput/TTFT/TPOT only over the saturated plateau to cut run-to-run variance at high concurrency; `steady_state` is an `Option` in JSON (`#[serde(default)]` for backward compat), null when the scope gate fails or `--no-steady-state`
+- **`--prompt-token-ids`** (random dataset only): send token-ID arrays instead of text to skip server-side tokenization; also skips the token-length verification pass (counts exact by construction)
+- **`--random-range-ratio`** follows Python semantics: lengths sampled uniformly from `[len*(1-r), len*(1+r)]`, default `0.0` = fixed; accepts a float in `[0,1)` or `'{"input": r1, "output": r2}'`. (The pre-2026-07 Rust-only form `[len*r, len]` with default 1.0 is rejected with a migration hint.)
+- **`prompt_list`** (`Arc<[Arc<str>]>` on `SampleRequest`/`RequestFuncInput`): multiple inputs per request for pooling backends — embeddings batches (`--random-batch-size`) send `"input": [...]`, rerank sends `[0]` as query + `[1..]` as documents
+- JSON output schema must match Python `vllm bench serve` exactly
+
+## Common Issues
+
+- **localhost vs 127.0.0.1**: Some systems resolve `localhost` to IPv6 `::1` while vLLM listens on IPv4 only. Use `127.0.0.1` or the actual hostname.
+- **Models without tokenizer.json** (e.g., `nvidia/Kimi-K2.5-NVFP4`): Automatically falls back to server-side tokenization. Can also use `--tokenizer` to point to a model with `tokenizer.json`.
+- **usage.completion_tokens parsing**: vLLM sends final usage chunk with `"choices":[]` (empty array). The usage `if` must be separate from the choices `if` (not `else if`).
+
+## Typical Usage
+
+```bash
+# Embedding benchmark (openai-embeddings, 8 inputs batched per request)
+./target/release/vllm-bench \
+  --backend openai-embeddings \
+  --base-url http://gb200-10:30000 \
+  --model BAAI/bge-large-en-v1.5 \
+  --dataset-name random \
+  --random-input-len 512 \
+  --random-batch-size 8 \
+  --num-prompts 1000 \
+  --save-result
+
+# vLLM rerank benchmark (one query + 8 documents per request)
+./target/release/vllm-bench \
+  --backend vllm-rerank \
+  --base-url http://gb200-10:30000 \
+  --model BAAI/bge-reranker-v2-m3 \
+  --dataset-name random-rerank \
+  --random-input-len 512 \
+  --random-batch-size 8 \
+  --num-prompts 500 \
+  --save-result
+
+# Prefix-cache stress (10 shared prefixes, 256+256 tokens)
+./target/release/vllm-bench \
+  --backend vllm \
+  --base-url http://gb200-10:30000 \
+  --model nvidia/Kimi-K2.5-NVFP4 \
+  --dataset-name prefix_repetition \
+  --prefix-repetition-prefix-len 256 \
+  --prefix-repetition-suffix-len 256 \
+  --prefix-repetition-num-prefixes 10 \
+  --num-prompts 1000
+
+# Custom JSONL workload ({"prompt": ..., "output_tokens": ...} per line)
+./target/release/vllm-bench \
+  --backend openai-chat \
+  --base-url http://gb200-10:30000 \
+  --model nvidia/Kimi-K2.5-NVFP4 \
+  --dataset-name custom \
+  --dataset-path workload.jsonl \
+  --custom-output-len -1 \
+  --num-prompts 1000
+
+# Random dataset
+./target/release/vllm-bench \
+  --backend vllm \
+  --base-url http://gb200-10:30000 \
+  --model nvidia/Kimi-K2.5-NVFP4 \
+  --dataset-name random \
+  --random-input-len 8192 \
+  --random-output-len 1024 \
+  --ignore-eos \
+  --num-prompts 4096 \
+  --percentile-metrics "ttft,tpot,itl,e2el" \
+  --save-result \
+  --max-concurrency 1400
+
+# Random multimodal dataset (VLM benchmark)
+./target/release/vllm-bench \
+  --backend openai-chat \
+  --base-url http://gb200-10:30000 \
+  --model Qwen/Qwen2.5-VL-7B-Instruct \
+  --dataset-name random-mm \
+  --random-input-len 512 \
+  --random-output-len 128 \
+  --num-prompts 100 \
+  --random-mm-base-items-per-request 1 \
+  --random-mm-limit-mm-per-prompt '{"image": 1, "video": 0}' \
+  --random-mm-bucket-config '{(1024, 800, 1): 1.0}'
+
+# HuggingFace dataset (WildChat)
+./target/release/vllm-bench \
+  --backend openai-chat \
+  --base-url http://gb200-10:30000 \
+  --model nvidia/Kimi-K2.5-NVFP4 \
+  --dataset-name hf \
+  --dataset-path allenai/WildChat-4.8M \
+  --hf-split train \
+  --num-prompts 1000 \
+  --save-result
+
+# HuggingFace dataset (LongBench with subset)
+./target/release/vllm-bench \
+  --backend openai-chat \
+  --base-url http://gb200-10:30000 \
+  --model nvidia/Kimi-K2.5-NVFP4 \
+  --dataset-name hf \
+  --dataset-path THUDM/LongBench \
+  --hf-subset narrativeqa \
+  --hf-split test \
+  --hf-output-len 512 \
+  --num-prompts 200
+```
@@ -0,0 +1 @@
+@AGENTS.md
@@ -0,0 +1,37 @@
+[package]
+name = "vllm-bench"
+version.workspace = true
+edition.workspace = true
+description = "High-performance benchmark client for vLLM serving endpoints"
+license.workspace = true
+
+[dependencies]
+anyhow.workspace = true
+base64.workspace = true
+bytes.workspace = true
+chrono.workspace = true
+clap.workspace = true
+dirs.workspace = true
+futures.workspace = true
+hf-hub.workspace = true
+image.workspace = true
+indicatif.workspace = true
+mimalloc.workspace = true
+rand.workspace = true
+rand_distr.workspace = true
+rayon.workspace = true
+reqwest = { workspace = true, features = ["json", "stream", "blocking", "http2"] }
+rlimit.workspace = true
+rustc-hash.workspace = true
+serde = { workspace = true, features = ["rc"] }
+serde_json = { workspace = true, features = ["raw_value"] }
+thiserror.workspace = true
+tiktoken-rs.workspace = true
+tokenizers.workspace = true
+tokio.workspace = true
+tokio-stream.workspace = true
+url.workspace = true
+uuid.workspace = true
+
+[lints]
+workspace = true
@@ -0,0 +1,810 @@
+# vllm-bench
+
+High-performance Rust benchmark client for vLLM serving endpoints. A drop-in replacement for `vllm bench serve` with near-instant startup, parallel dataset generation, and a fraction of the memory overhead — and no Python at runtime.
+
+```bash
+vllm-bench --backend vllm --base-url http://127.0.0.1:8000 \
+  --model <model> --dataset-name random \
+  --random-input-len 1024 --random-output-len 128 \
+  --num-prompts 1000 --max-concurrency 200
+```
+
+## Highlights
+
+- **Fast** — ~7 ms startup, single ~7 MB static binary, no Python imports.
+- **Scales** — `Arc<str>` prompt sharing + mimalloc keep memory <100 MB at 1400+ concurrency.
+- **Many datasets** — `random`, `random-mm` (VLM), `sharegpt`, `sonnet`, `speed-bench`, and any HuggingFace dataset.
+- **Many backends** — completions, chat, embeddings, pooling, and rerank.
+- **Beyond a single run** — concurrency/rate **sweeps**, **multi-run** stats, **multi-turn** conversations, **LoRA** multi-adapter, and result **comparison**.
+- **Steady-state metrics** — throughput/latency measured over the saturated plateau, excluding ramp-up and drain.
+- **Parity** — JSON output schema and timing semantics match Python `vllm bench serve` exactly.
+
+### Performance vs. Python
+
+| Metric | Python | Rust |
+| -------- | -------- | ------ |
+| Startup time | Multi-second (import vllm + numpy + aiohttp) | ~7 ms |
+| 100k random prompts (input_len=8192) | Minutes | Seconds (rayon parallelism) |
+| Binary size | — | ~7 MB |
+| Peak memory at 1400 concurrency | High (GIL + per-object overhead) | <100 MB (`Arc<str>` prompt sharing) |
+
+## Contents
+
+- [Install](#install)
+- [Quick Start](#quick-start)
+- [Usage Examples](#usage-examples)
+- [Supported Backends](#supported-backends)
+- [Supported Datasets](#supported-datasets)
+- [Metrics](#metrics)
+- [CLI Reference](#cli-reference)
+- [Tokenizer Support](#tokenizer-support)
+- [Output Format](#output-format)
+- [Architecture](#architecture)
+- [Environment Variables](#environment-variables)
+
+## Install
+
+### Prebuilt binaries (Linux)
+
+```bash
+curl -fsSL https://github.com/vllm-project/vllm-bench/releases/latest/download/vllm-bench-$(uname -m)-linux-musl -o vllm-bench && chmod +x vllm-bench
+```
+
+### With Cargo
+
+Install straight from the repository (builds from source; requires [Rust](https://rustup.rs/) stable and a C compiler for the native tokenizer dependency):
+
+```bash
+cargo install --git https://github.com/vllm-project/vllm-bench vllm-bench
+```
+
+The trailing `vllm-bench` selects the package — the repo also ships a `mock-llm-server` binary, so omitting it fails with `multiple packages with binaries found`. The binary is installed to `~/.cargo/bin/`.
+
+### Build from source
+
+Requires [Rust](https://rustup.rs/) (stable).
+
+```bash
+git clone https://github.com/vllm-project/vllm-bench.git
+cd vllm-bench
+./install.sh            # builds release and installs to ~/.local/bin
+# or: ./install.sh --to ~/bin
+```
+
+## Quick Start
+
+Point it at a running vLLM server and benchmark with synthetic prompts:
+
+```bash
+vllm-bench \
+  --backend vllm \
+  --base-url http://127.0.0.1:8000 \
+  --model <model-name> \
+  --dataset-name random \
+  --random-input-len 1024 \
+  --random-output-len 128 \
+  --num-prompts 1000 \
+  --max-concurrency 200
+```
+
+> **Tip:** prefer `127.0.0.1` over `localhost` — some systems resolve `localhost` to IPv6 `::1` while vLLM listens on IPv4 only.
+
+Add `--save-result` to write a JSON file, or `--dry-run` to generate and inspect the dataset without sending any requests.
+
+## Usage Examples
+
+<details open>
+<summary><b>Generation (completions / chat)</b></summary>
+
+```bash
+# Full production-style run with percentile metrics and result file
+vllm-bench \
+  --backend vllm \
+  --base-url http://127.0.0.1:8000 \
+  --model nvidia/Kimi-K2.5-NVFP4 \
+  --dataset-name random \
+  --random-input-len 8192 \
+  --random-output-len 1024 \
+  --ignore-eos \
+  --num-prompts 4096 \
+  --percentile-metrics "ttft,tpot,itl,e2el" \
+  --save-result \
+  --max-concurrency 1400
+
+# Send token IDs instead of text (pure vLLM: skips server-side tokenization,
+# exact token counts, faster). Random dataset only.
+vllm-bench \
+  --backend vllm \
+  --base-url http://127.0.0.1:8000 \
+  --model <model-name> \
+  --dataset-name random \
+  --random-input-len 1024 \
+  --prompt-token-ids \
+  --num-prompts 1000
+```
+
+</details>
+
+<details>
+<summary><b>Datasets (ShareGPT / Sonnet / HuggingFace / SPEED-Bench)</b></summary>
+
+```bash
+# ShareGPT (auto-downloads from HuggingFace on first run, cached afterwards)
+vllm-bench \
+  --backend openai-chat --base-url http://127.0.0.1:8000 --model <model-name> \
+  --dataset-name sharegpt --num-prompts 500 --save-result
+
+# ShareGPT with an explicit local file
+vllm-bench \
+  --backend openai-chat --base-url http://127.0.0.1:8000 --model <model-name> \
+  --dataset-name sharegpt --dataset-path /path/to/ShareGPT_V3.json \
+  --num-prompts 500 --save-result
+
+# Sonnet — built-in Shakespeare sonnets, no dataset file needed.
+# Generates prompts of a controllable token length with a shared prefix.
+vllm-bench \
+  --backend openai-chat --base-url http://127.0.0.1:8000 --model <model-name> \
+  --dataset-name sonnet \
+  --sonnet-input-len 550 --sonnet-output-len 150 --sonnet-prefix-len 200 \
+  --num-prompts 500
+
+# Any public HuggingFace dataset (auto-downloads, auto-detects columns)
+vllm-bench \
+  --backend openai-chat --base-url http://127.0.0.1:8000 --model <model-name> \
+  --dataset-name hf --dataset-path allenai/WildChat-4.8M \
+  --hf-split train --num-prompts 1000 --save-result
+
+# HuggingFace dataset with subset + fixed output length (LongBench)
+vllm-bench \
+  --backend openai-chat --base-url http://127.0.0.1:8000 --model <model-name> \
+  --dataset-name hf --dataset-path THUDM/LongBench \
+  --hf-subset narrativeqa --hf-split test --hf-output-len 512 --num-prompts 200
+
+# Gated HuggingFace dataset (requires HF_TOKEN)
+HF_TOKEN=hf_xxx vllm-bench \
+  --backend openai-chat --base-url http://127.0.0.1:8000 --model <model-name> \
+  --dataset-name hf --dataset-path lmsys/lmsys-chat-1m \
+  --hf-split train --hf-output-len 256 --num-prompts 1000
+
+# SPEED-Bench for speculative decoding evaluation (auto-downloads, cached)
+vllm-bench \
+  --backend openai-chat --base-url http://127.0.0.1:8000 --model <model-name> \
+  --dataset-name speed-bench --speed-bench-config qualitative \
+  --num-prompts 200 --output-len 256 --save-result
+
+# SPEED-Bench throughput split with entropy category filter + input truncation
+vllm-bench \
+  --backend openai-chat --base-url http://127.0.0.1:8000 --model <model-name> \
+  --dataset-name speed-bench --speed-bench-config throughput_16k \
+  --speed-bench-max-input-len 10240 --speed-bench-category low_entropy \
+  --num-prompts 500 --output-len 256 --max-concurrency 200 --save-result
+```
+
+</details>
+
+<details>
+<summary><b>Multimodal (VLM with synthetic images)</b></summary>
+
+```bash
+# One synthetic image per request
+vllm-bench \
+  --backend openai-chat --base-url http://127.0.0.1:8000 \
+  --model Qwen/Qwen2.5-VL-7B-Instruct \
+  --dataset-name random-mm \
+  --random-input-len 512 --random-output-len 128 --num-prompts 100 \
+  --random-mm-base-items-per-request 1 \
+  --random-mm-limit-mm-per-prompt '{"image": 1, "video": 0}' \
+  --random-mm-bucket-config '{(1024, 800, 1): 1.0}'
+
+# Multiple images per request, mixed resolutions
+vllm-bench \
+  --backend openai-chat --base-url http://127.0.0.1:8000 \
+  --model Qwen/Qwen2.5-VL-7B-Instruct \
+  --dataset-name random-mm \
+  --random-input-len 256 --random-output-len 128 --num-prompts 50 \
+  --random-mm-base-items-per-request 3 \
+  --random-mm-limit-mm-per-prompt '{"image": 5, "video": 0}' \
+  --random-mm-bucket-config '{(256,256,1): 0.5, (720,1280,1): 0.5}'
+```
+
+</details>
+
+<details>
+<summary><b>Embedding / Pooling / Rerank</b></summary>
+
+```bash
+# Text embedding
+vllm-bench \
+  --backend openai-embeddings --base-url http://127.0.0.1:8000 \
+  --model BAAI/bge-large-en-v1.5 \
+  --dataset-name random --random-input-len 512 --num-prompts 1000 \
+  --max-concurrency 200 --save-result
+
+# Chat-format embedding (supports multimodal content)
+vllm-bench \
+  --backend openai-embeddings-chat --base-url http://127.0.0.1:8000 \
+  --model BAAI/bge-large-en-v1.5 \
+  --dataset-name sharegpt --num-prompts 500 --save-result
+
+# vLLM native pooling endpoint
+vllm-bench \
+  --backend vllm-pooling --base-url http://127.0.0.1:8000 \
+  --model BAAI/bge-large-en-v1.5 \
+  --dataset-name random --random-input-len 256 --num-prompts 1000 --save-result
+
+# Rerank (query from dataset, documents via --extra-body)
+vllm-bench \
+  --backend vllm-rerank --base-url http://127.0.0.1:8000 \
+  --model BAAI/bge-reranker-v2-m3 \
+  --dataset-name sharegpt --num-prompts 500 \
+  --extra-body '{"documents": ["document to rerank"]}' --save-result
+```
+
+</details>
+
+<details>
+<summary><b>Rate control, ramp-up &amp; goodput</b></summary>
+
+```bash
+# Ramp from 10 → 100 RPS with goodput SLO tracking
+vllm-bench \
+  --backend vllm --base-url http://127.0.0.1:8000 --model <model-name> \
+  --num-prompts 2000 \
+  --ramp-up-strategy linear --ramp-up-start-rps 10 --ramp-up-end-rps 100 \
+  --goodput ttft:200 e2el:5000 \
+  --save-result
+
+# Fixed Poisson arrival rate at 50 RPS
+vllm-bench \
+  --backend vllm --base-url http://127.0.0.1:8000 --model <model-name> \
+  --num-prompts 2000 --request-rate 50 --burstiness 1.0
+```
+
+</details>
+
+<details>
+<summary><b>Sweep — find the optimal concurrency / rate</b></summary>
+
+```bash
+# Sweep over concurrency values
+vllm-bench \
+  --backend vllm --base-url http://127.0.0.1:8000 --model <model-name> \
+  --num-prompts 500 \
+  --sweep-max-concurrency 1,10,50,100,200,500,1000
+
+# Sweep over request rates
+vllm-bench \
+  --backend vllm --base-url http://127.0.0.1:8000 --model <model-name> \
+  --num-prompts 500 \
+  --sweep-request-rate 1,10,50,100,inf
+
+# Scale work with concurrency and reset the prefix cache between points
+# (--sweep-num-prompts-factor sets num_prompts = concurrency * factor;
+#  --reset-prefix-cache requires VLLM_SERVER_DEV_MODE=1 on the server)
+vllm-bench \
+  --backend vllm --base-url http://127.0.0.1:8000 --model <model-name> \
+  --sweep-max-concurrency 1,10,50,100 \
+  --sweep-num-prompts-factor 20 \
+  --reset-prefix-cache
+```
+
+</details>
+
+<details>
+<summary><b>Multi-run &amp; comparison</b></summary>
+
+```bash
+# Run 5 times, report mean/std/min/max with coefficient of variation
+vllm-bench \
+  --backend vllm --base-url http://127.0.0.1:8000 --model <model-name> \
+  --num-prompts 1000 --max-concurrency 200 --num-runs 5
+
+# Compare two saved result files side-by-side (no server needed)
+vllm-bench --compare baseline.json optimized.json
+```
+
+</details>
+
+<details>
+<summary><b>Multi-turn conversations</b></summary>
+
+```bash
+# Synthetic multi-turn (controllable per-turn token lengths)
+vllm-bench \
+  --backend openai-chat --base-url http://127.0.0.1:8000 --model <model-name> \
+  --dataset-name random --multi-turn --multi-turn-num-turns 5 \
+  --random-input-len 512 --random-output-len 256 \
+  --num-prompts 50 --multi-turn-concurrency 10 \
+  --percentile-metrics "ttft,tpot,itl,e2el" --save-result
+
+# Variable turn count per conversation + per-turn input length for turns 1+
+vllm-bench \
+  --backend openai-chat --base-url http://127.0.0.1:8000 --model <model-name> \
+  --dataset-name random --multi-turn \
+  --multi-turn-min-turns 2 --multi-turn-max-turns 8 \
+  --random-input-len 2048 --per-turn-input-len 256 --random-output-len 128 \
+  --num-prompts 100 --multi-turn-concurrency 20
+
+# ShareGPT conversations (loads all turns, not just the first two)
+vllm-bench \
+  --backend openai-chat --base-url http://127.0.0.1:8000 --model <model-name> \
+  --dataset-name sharegpt --multi-turn \
+  --num-prompts 50 --multi-turn-concurrency 10 --save-result
+
+# Think time between turns
+vllm-bench \
+  --backend openai-chat --base-url http://127.0.0.1:8000 --model <model-name> \
+  --multi-turn --multi-turn-num-turns 3 --multi-turn-delay-ms 500 \
+  --num-prompts 100 --multi-turn-concurrency 20
+```
+
+</details>
+
+<details>
+<summary><b>LoRA multi-adapter</b></summary>
+
+```bash
+# Distribute requests across N adapters registered on the server.
+# --model stays the BASE model (tokenizer / readiness / /tokenize use it);
+# the per-request `model` field is rewritten to one of --lora-modules.
+vllm-bench \
+  --backend openai-chat --base-url http://127.0.0.1:8000 \
+  --model Qwen/Qwen3-30B-A3B \
+  --lora-modules sql-lora-1 sql-lora-2 sql-lora-3 sql-lora-4 \
+                 sql-lora-5 sql-lora-6 sql-lora-7 sql-lora-8 \
+  --lora-assignment random \
+  --dataset-name random --random-input-len 1024 --random-output-len 256 \
+  --num-prompts 1000 --max-concurrency 64 --save-result
+
+# Deterministic round-robin assignment (request i -> adapter[i % N])
+vllm-bench \
+  --backend openai-chat --base-url http://127.0.0.1:8000 \
+  --model Qwen/Qwen3-30B-A3B \
+  --lora-modules sql-lora-1 sql-lora-2 sql-lora-3 sql-lora-4 \
+  --lora-assignment round-robin \
+  --dataset-name random --num-prompts 1000
+```
+
+Server side — start vLLM with `--enable-lora` and one `name=path` pair per adapter:
+
+```bash
+vllm serve <base-model> \
+  --enable-lora --max-loras 8 --max-lora-rank 16 \
+  --lora-modules \
+    sql-lora-1=jeeejeee/qwen3-moe-text2sql-spider \
+    sql-lora-2=jeeejeee/qwen3-moe-text2sql-spider \
+    ...
+```
+
+Set `--max-loras` ≥ number of adapter names to keep them all resident (clean steady-state numbers), or lower to stress the LoRA swap path.
+
+</details>
+
+<details>
+<summary><b>Profiling &amp; dry-run</b></summary>
+
+```bash
+# Trigger vLLM server-side profiling (start before, stop after the benchmark)
+vllm-bench \
+  --backend vllm --base-url http://127.0.0.1:8000 --model <model-name> \
+  --num-prompts 100 --profile
+
+# Defer profiling until the server batch is full, then capture for 10s
+vllm-bench \
+  --backend vllm --base-url http://127.0.0.1:8000 --model <model-name> \
+  --num-prompts 2000 --max-concurrency 256 \
+  --profile --profile-batch-threshold 200 --profile-duration 10
+
+# Dry run: generate dataset, print stats, send nothing
+vllm-bench \
+  --model <model-name> --num-prompts 100000 --random-input-len 8192 --dry-run
+```
+
+</details>
+
+## Supported Backends
+
+### Generation
+
+| Backend | API Endpoint | Description |
+| --------- | ------------- | ------------- |
+| `vllm` / `openai` | `/v1/completions` | OpenAI-compatible completions (streaming) |
+| `openai-chat` | `/v1/chat/completions` | OpenAI-compatible chat completions (streaming, multimodal) |
+
+### Embedding / Pooling
+
+| Backend | API Endpoint | Description |
+| --------- | ------------- | ------------- |
+| `openai-embeddings` | `/v1/embeddings` | Text embedding (accepts text or token IDs) |
+| `openai-embeddings-chat` | `/v1/embeddings` | Chat-format embedding (supports multimodal content) |
+| `vllm-pooling` | `/v1/pooling` | vLLM native pooling endpoint |
+| `vllm-rerank` | `/v1/rerank` | vLLM reranking (query from prompt, documents via `--extra-body`) |
+
+Pooling backends are non-streaming and report E2EL (end-to-end latency) only. Use `--dataset-name sharegpt`, `sonnet`, or `hf` for text-based embedding/rerank benchmarks, or `random` for token-ID-based embedding benchmarks.
+
+## Supported Datasets
+
+| Dataset | Description |
+| --------- | ------------- |
+| `random` | Synthetic prompts with exact token-length matching (default) |
+| `random-mm` | Synthetic multimodal prompts with random JPEG images for VLM benchmarking (requires `openai-chat`) |
+| `sharegpt` | Real conversations from ShareGPT (auto-downloads from HuggingFace, or use `--dataset-path`) |
+| `sonnet` | Built-in Shakespeare sonnets; controllable token length + shared prefix, no dataset file needed |
+| `speed-bench` | NVIDIA SPEED-Bench for speculative decoding evaluation (auto-downloads, 11 categories) |
+| `hf` | Any HuggingFace dataset (auto-downloads via datasets-server API, auto-detects chat/text columns) |
+
+## Metrics
+
+### Generation backends
+
+- **TTFT** (Time to First Token) — latency from request send to first token received
+- **TPOT** (Time per Output Token) — average time between output tokens
+- **ITL** (Inter-Token Latency) — per-token latency distribution
+- **E2EL** (End-to-End Latency) — total request latency
+- **Throughput** — requests/sec, output tokens/sec, peak output tokens/sec, total tokens/sec
+- **Concurrency** — peak concurrent requests
+- **Goodput** — requests/sec meeting all specified SLOs (with `--goodput`)
+
+### Pooling / embedding backends
+
+- **E2EL** — total request latency (mean, median, std, percentiles)
+- **Throughput** — requests/sec, input tokens/sec
+- **Concurrency** — peak concurrent requests
+
+### Steady-state metrics
+
+When `--max-concurrency` is set and `--request-rate` is `inf` (closed-loop mode), the benchmark automatically reports an additional **Steady-State Metrics** block. It measures throughput and latency only over the window during which in-flight concurrency stays at or above a fraction of `--max-concurrency`, excluding the ramp-up and drain phases. This sharply reduces run-to-run variance at very high concurrency.
+
+The block reports request/input/output/total token throughput plus TTFT (mean, median, percentiles) and TPOT (mean, median, P90, P99) over the detected plateau, along with the window bounds and how many requests fell inside it. Tune it with `--steady-state-threshold` (default `0.95`) and `--steady-state-min-window`, or disable with `--no-steady-state`. The result JSON carries a `steady_state` object (null when not computed).
+
+## CLI Reference
+
+Run `vllm-bench --help` for the authoritative list. Grouped reference below.
+
+<details>
+<summary><b>Server connection</b></summary>
+
+| Flag | Default | Description |
+| ------ | --------- | ------------- |
+| `--backend` | `openai` | Backend type (`vllm`, `openai`, `openai-chat`, `openai-embeddings`, `openai-embeddings-chat`, `vllm-pooling`, `vllm-rerank`) |
+| `--base-url` | — | Server base URL (overrides `--host`/`--port`) |
+| `--host` | `127.0.0.1` | Server host |
+| `--port` | `8000` | Server port |
+| `--endpoint` | Auto | API endpoint path (auto-selected per backend) |
+| `--insecure` | `false` | Disable SSL certificate verification |
+
+</details>
+
+<details>
+<summary><b>Model &amp; tokenizer</b></summary>
+
+| Flag | Default | Description |
+| ------ | --------- | ------------- |
+| `--model` | Auto-detect | Model name (fetched from `/v1/models` if omitted) |
+| `--served-model-name` | — | Model name used in API requests |
+| `--tokenizer` | Same as model | Tokenizer name or path (supports HF, tiktoken, server fallback) |
+| `--tokenizer-mode` | `auto` | Tokenizer mode (`auto`, `hf`, `slow`, `mistral`) |
+| `--trust-remote-code` | `false` | Trust remote code for tokenizer |
+| `--skip-tokenizer-init` | `false` | Skip tokenizer initialization |
+
+</details>
+
+<details>
+<summary><b>Dataset</b></summary>
+
+| Flag | Default | Description |
+| ------ | --------- | ------------- |
+| `--dataset-name` | `random` | Dataset type (`random`, `random-mm`, `sharegpt`, `sonnet`, `speed-bench`, `hf`) |
+| `--dataset-path` | — | Path to dataset file (optional for `sharegpt`/`sonnet`, which auto-source) |
+| `--num-prompts` | `1000` | Number of prompts to generate (conversations in multi-turn mode) |
+| `--max-model-len` | — | Filter out requests where `prompt_len + output_len` exceeds this context length |
+| `--input-len` | — | Override input length (general) |
+| `--output-len` | — | Override output length (general) |
+| `--no-oversample` | `false` | Don't oversample if dataset is smaller than `--num-prompts` |
+| `--disable-shuffle` | `false` | Don't shuffle the dataset |
+| `--seed` | `0` | Random seed for reproducibility |
+| **Random** | | |
+| `--random-input-len` | `1024` | Input token length |
+| `--random-output-len` | `128` | Output token length |
+| `--random-prefix-len` | `0` | Shared prefix length |
+| `--random-range-ratio` | `1.0` | Length jitter, range `(0, 1]`. Lengths sampled from `[ratio × target, target]`; `1.0` = fixed length |
+| `--prompt-token-ids` | `false` | Send prompts as token-ID arrays (skips server-side tokenization, exact counts). Random dataset only |
+| **Random multimodal** | | |
+| `--random-mm-base-items-per-request` | `1` | Base number of multimodal items (images) per request |
+| `--random-mm-num-mm-items-range-ratio` | `0.0` | Range ratio for varying item count per request |
+| `--random-mm-limit-mm-per-prompt` | `{"image": 255, "video": 1}` | Per-modality hard caps (JSON) |
+| `--random-mm-bucket-config` | `{(256,256,1): 0.5, (720,1280,1): 0.5}` | `(height,width,frames)` → probability (Python tuple syntax; frames=1 = image) |
+| **ShareGPT** | | |
+| `--sharegpt-output-len` | — | Override output length |
+| **Sonnet** | | |
+| `--sonnet-input-len` | `550` | Input tokens per request |
+| `--sonnet-output-len` | `150` | Output tokens per request |
+| `--sonnet-prefix-len` | `200` | Prefix tokens shared across requests |
+| **SPEED-Bench** | | |
+| `--speed-bench-config` | `qualitative` | Split (`qualitative`, `throughput_1k`/`2k`/`8k`/`16k`/`32k`) |
+| `--speed-bench-category` | — | Filter by category (`low_entropy`, `high_entropy`, `mixed_entropy`, `coding`, `math`, …) |
+| `--speed-bench-max-input-len` | — | Truncate prompts to at most N tokens |
+| **HuggingFace** | | |
+| `--hf-split` | Auto | Split (`train`, `test`, `validation`); auto-detected if omitted |
+| `--hf-subset` | — | Subset/config name (e.g. `narrativeqa` for LongBench) |
+| `--hf-output-len` | — | Fixed output length for all requests (overrides dataset-derived length) |
+| `--hf-text-column` | Auto | Column containing prompt text; auto-detected from common patterns |
+
+</details>
+
+<details>
+<summary><b>Rate control</b></summary>
+
+| Flag | Default | Description |
+| ------ | --------- | ------------- |
+| `--request-rate` | `inf` | Requests per second (`inf` = all at once) |
+| `--burstiness` | `1.0` | Burstiness factor (1.0 = Poisson, >1 = bursty) |
+| `--max-concurrency` | `num-prompts` | Maximum concurrent requests (semaphore) |
+| `--ramp-up-strategy` | — | Ramp-up mode (`linear` or `exponential`) |
+| `--ramp-up-start-rps` | — | Starting request rate for ramp-up |
+| `--ramp-up-end-rps` | — | Ending request rate for ramp-up |
+
+</details>
+
+<details>
+<summary><b>Sampling parameters</b></summary>
+
+| Flag | Description |
+| ------ | ------------- |
+| `--temperature` | Temperature (server default if omitted) |
+| `--top-p` | Top-p (nucleus) sampling |
+| `--top-k` | Top-k sampling |
+| `--min-p` | Min-p sampling |
+| `--frequency-penalty` | Frequency penalty |
+| `--presence-penalty` | Presence penalty |
+| `--repetition-penalty` | Repetition penalty |
+
+Merged into the request body. Only effective with generation backends (`vllm`, `openai`, `openai-chat`); ignored by pooling/embedding backends.
+
+</details>
+
+<details>
+<summary><b>Output &amp; results</b></summary>
+
+| Flag | Default | Description |
+| ------ | --------- | ------------- |
+| `--save-result` | `false` | Save results to JSON file |
+| `--save-detailed` | `false` | Include per-request data in JSON (input/output lens, ITLs, texts) |
+| `--append-result` | `false` | Append to existing JSON file (JSONL format) |
+| `--result-dir` | — | Directory for result files |
+| `--result-filename` | Auto | Custom result filename |
+| `--percentile-metrics` | `ttft,tpot,itl,e2el` | Metrics for percentile reporting (pooling defaults to `e2el` only) |
+| `--metric-percentiles` | `99` | Percentile values to compute |
+| `--sweep-summary-percentiles` | — | Extra percentiles for sweep summary tables (auto-added to computed set) |
+| `--goodput` | — | SLO pairs for goodput (`ttft:100 tpot:50 e2el:500`, values in ms) |
+| `--disable-tqdm` | `false` | Disable progress bar |
+| `--label` | — | Label prefix for result files |
+| `--metadata` | — | Key-value metadata (`KEY=VALUE`, repeatable) |
+
+</details>
+
+<details>
+<summary><b>Request options</b></summary>
+
+| Flag | Default | Description |
+| ------ | --------- | ------------- |
+| `--ignore-eos` | `false` | Ignore EOS token (force full output length) |
+| `--logprobs` | — | Number of logprobs per token |
+| `--num-warmups` | `0` | Warmup requests before benchmarking |
+| `--ready-check-timeout-sec` | `0` | Endpoint readiness timeout (0 = skip) |
+| `--request-id-prefix` | Auto (UUID) | Prefix for request IDs |
+| `--header` | — | Extra headers (`KEY=VALUE`, repeatable) |
+| `--extra-body` | — | Extra JSON body parameters |
+| `--dry-run` | `false` | Generate dataset only, skip benchmark |
+
+</details>
+
+<details>
+<summary><b>Steady-state metrics</b></summary>
+
+| Flag | Default | Description |
+| ------ | --------- | ------------- |
+| `--steady-state-threshold` | `0.95` | Fraction of `--max-concurrency` at which the steady-state window opens, range (0, 1] |
+| `--steady-state-min-window` | Auto | Minimum window duration (s) below which a warning is attached. Default `max(10, 0.1 × run_duration)` |
+| `--no-steady-state` | `false` | Disable steady-state metrics computation |
+
+Computed only when `--max-concurrency` is set and `--request-rate` is `inf`.
+
+</details>
+
+<details>
+<summary><b>Profiling</b></summary>
+
+| Flag | Default | Description |
+| ------ | --------- | ------------- |
+| `--profile` | `false` | Trigger vLLM server-side profiling (`/start_profile` before, `/stop_profile` after) |
+| `--profile-batch-threshold` | — | Defer profiling until `/metrics` reports ≥ N running requests, then capture. Requires `--profile` |
+| `--profile-duration` | `5.0` | Seconds to capture once the batch threshold is reached. Requires `--profile-batch-threshold` |
+
+</details>
+
+<details>
+<summary><b>Sweep mode</b></summary>
+
+| Flag | Default | Description |
+| ------ | --------- | ------------- |
+| `--sweep-max-concurrency` | — | Comma-separated concurrency values to sweep (e.g. `1,10,50,100,500`) |
+| `--sweep-request-rate` | — | Comma-separated rate values to sweep, supports `inf` (e.g. `1,10,100,inf`) |
+| `--sweep-num-prompts-factor` | — | Set `num_prompts = concurrency × factor` per concurrency sweep point |
+| `--reset-prefix-cache` | `false` | Reset the server's prefix cache before each sweep iteration (requires `VLLM_SERVER_DEV_MODE=1`) |
+
+Runs the benchmark once per value, then prints a summary table comparing throughput and latency across all sweep points and identifies the best-throughput configuration. Works in multi-turn mode too. `--sweep-summary-percentiles` appends extra TTFT/TPOT/E2EL columns to the summary, auto-adding any missing percentiles to the computed set so they also appear in result JSON.
+
+</details>
+
+<details>
+<summary><b>Multi-turn conversation benchmark</b></summary>
+
+| Flag | Default | Description |
+| ------ | --------- | ------------- |
+| `--multi-turn` | `false` | Enable multi-turn conversation mode (requires `--backend openai-chat`) |
+| `--multi-turn-num-turns` | `3` | Turns per conversation (synthetic mode) |
+| `--multi-turn-min-turns` | `0` | Minimum turns per conversation (0 = use `--multi-turn-num-turns`) |
+| `--multi-turn-max-turns` | `0` | Maximum turns per conversation (0 = `--multi-turn-num-turns` synthetic / uncapped ShareGPT) |
+| `--multi-turn-concurrency` | — | Concurrent conversations (defaults to `--max-concurrency` or `--num-prompts`) |
+| `--multi-turn-delay-ms` | `0` | Delay between turns in ms (simulates user think time) |
+| `--per-turn-input-len` | `0` | Input token length for turns 1+ (0 = use `--random-input-len` for all turns) |
+| `--multi-turn-prefix-global-ratio` | `0.0` | Fraction of per-turn input shared across all conversations (random dataset only) |
+| `--multi-turn-prefix-conversation-ratio` | `0.0` | Fraction shared within each conversation (random dataset only) |
+
+With `--multi-turn`, `--num-prompts` controls the number of **conversations**, not individual requests.
+
+**How it works:**
+
+- Turn 1: send `[user_1]`, get `assistant_1`
+- Turn 2: send `[user_1, assistant_1, user_2]`, get `assistant_2`
+- Turn N: send full history + `user_N` — measures growing-context performance
+
+**Data sources:**
+
+- `--dataset-name random` — synthetic conversations with controllable per-turn token lengths. Auto-sets `min_tokens` to enforce output length without `ignore_eos`.
+- `--dataset-name sharegpt` — loads all turns (not just the first two); filters for entries with ≥ 2 real turns.
+
+**Prefix sharing** (random dataset): when `--multi-turn-prefix-global-ratio` or `--multi-turn-prefix-conversation-ratio` is > 0, each turn sends a fixed-length message (no history accumulation) composed of a global prefix + per-conversation prefix + unique suffix. The two ratios must sum to < 1.0.
+
+**Router affinity:** every turn sends `X-Session-ID: {conversation_id}` for KV-cache reuse behind a vLLM router.
+
+**Output:** overall metrics plus a per-turn breakdown (TTFT/TPOT/ITL/E2EL by turn index). Expect TTFT to climb across turns due to growing context. JSON includes a `per_turn_metrics` array.
+
+</details>
+
+<details>
+<summary><b>LoRA multi-adapter</b></summary>
+
+| Flag | Default | Description |
+| ------ | --------- | ------------- |
+| `--lora-modules` | — | Adapter names registered on the server (`vllm serve --lora-modules name=path`). Each request's `model` field is rewritten to one of these. Repeatable |
+| `--lora-assignment` | `random` | Distribution: `random` (uniform, seeded by `--seed`) or `round-robin` (deterministic `i % N`) |
+
+`--model` must stay the **base** model — its tokenizer builds prompts, and `/v1/models`, `/tokenize`, ready check, and warmup all use it. Only the per-request `model` field in completions/chat payloads is rewritten to the assigned adapter (vLLM routes by name).
+
+**Assignment scope:** per request in single-shot mode; **per conversation** (sticky across all turns) in multi-turn mode, to avoid breaking prefix-cache reuse mid-dialog.
+
+**Reproducibility:** with `--lora-assignment random`, the same `--seed` + same `--lora-modules` list yields identical request-to-adapter mappings. Pooling/embedding backends are rejected — LoRA routing applies to generative paths only.
+
+</details>
+
+<details>
+<summary><b>Multi-run &amp; comparison</b></summary>
+
+| Flag | Default | Description |
+| ------ | --------- | ------------- |
+| `--num-runs` | `1` | Run benchmark N times; report mean/std/min/max with CV |
+| `--compare` | — | Compare two result JSON files side-by-side (skips benchmarking) |
+
+`--num-runs` aggregates metrics across runs and reports the coefficient of variation (CV) for throughput stability. `--compare` reads two previously-saved result files and prints a diff with delta, % change, and improvement/regression markers.
+
+</details>
+
+## Tokenizer Support
+
+Tokenizers are loaded with a three-tier fallback chain:
+
+1. **Local HuggingFace** — `tokenizer.json` from a local path or the Hub (fastest)
+2. **Tiktoken** — `.tiktoken` / `.model` format for Kimi, Qwen, etc. (auto-extracts `pat_str` from Python source)
+3. **Server-side** — falls back to vLLM's `/tokenize` + `/detokenize` endpoints
+
+For the `random` dataset, prompt token lengths are verified against the server on the first run and cached; subsequent runs with the same model+server skip verification. Verification is also skipped when `--prompt-token-ids` is set (token counts are exact by construction).
+
+Models without `tokenizer.json` (e.g. `nvidia/Kimi-K2.5-NVFP4`) fall back to server-side tokenization automatically; you can also point `--tokenizer` at a model that ships `tokenizer.json`.
+
+## Output Format
+
+JSON output is compatible with the `vllm bench serve` Python schema. Result files are named:
+
+```text
+{label}-{rate}qps-concurrency{max_concurrency}-{model}-{timestamp}.json
+```
+
+Use `--append-result` to append multiple runs to the same file in JSONL format. `--save-detailed` adds per-request arrays (input/output lengths, ITLs, generated text).
+
+## Architecture
+
+<details>
+<summary><b>Source layout</b></summary>
+
+```text
+src/
+├── main.rs                  # Entry point, mimalloc, tokio runtime, mode dispatch
+├── cli.rs                   # clap CLI argument definitions
+├── config.rs                # Validated config, goodput/ramp-up parsing
+├── benchmark.rs             # Core orchestrator (schedule, spawn, collect, verify, profile)
+├── multi_turn.rs            # Multi-turn conversation orchestrator (channel workers)
+├── compare.rs               # Result diff (--compare file_a.json file_b.json)
+├── sweep.rs                 # Parameter sweep (--sweep-max-concurrency, --sweep-request-rate)
+├── multi_run.rs             # Multi-run statistics (--num-runs N)
+├── rate_control.rs          # Gamma/Poisson scheduling + linear/exponential ramp-up
+├── ready_checker.rs         # Endpoint readiness with retry
+├── tokenizer.rs             # Tokenizer abstraction (HF, tiktoken, server)
+├── tiktoken.rs              # Tiktoken BPE loader with pat_str extraction
+├── error.rs                 # Error types
+├── backends/
+│   ├── mod.rs               # Backend enum dispatch, typed SSE structs
+│   ├── streaming.rs         # SSE stream parser with speculative JSON parse
+│   ├── openai_completions.rs # /v1/completions backend
+│   ├── openai_chat.rs       # /v1/chat/completions backend
+│   └── pooling.rs           # Embedding/pooling/rerank backends (non-streaming)
+├── datasets/
+│   ├── mod.rs               # SampleRequest, ConversationTurn, MultiTurnConversation types
+│   ├── random.rs            # Random dataset with rayon parallelism
+│   ├── random_mm.rs         # Random multimodal dataset (JPEG generation, bucket sampling)
+│   ├── multi_turn.rs        # Multi-turn synthetic + ShareGPT conversation generators
+│   ├── sharegpt.rs          # ShareGPT JSON dataset loader
+│   ├── sonnet.rs            # Sonnet dataset (built-in Shakespeare sonnets)
+│   ├── speed_bench.rs       # NVIDIA SPEED-Bench loader (auto-download + cache)
+│   └── hf_dataset.rs        # Generic HuggingFace dataset (auto-download, column detection)
+├── metrics/
+│   ├── mod.rs               # BenchmarkMetrics, MultiTurnMetrics structs
+│   ├── calculator.rs        # Percentile/throughput/goodput/peak/multi-turn computation
+│   └── steady_state.rs      # Steady-state window detection + plateau metrics
+└── output/
+    ├── mod.rs
+    ├── console.rs           # Terminal output (matches Python format)
+    └── json.rs              # JSON result serialization (Python-compatible schema)
+```
+
+</details>
+
+### Key design decisions
+
+- **reqwest + tokio** — HTTP client with connection pooling, forced HTTP/1.1, TCP_NODELAY to match Python's aiohttp and avoid Nagle latency inflation on TTFT
+- **mimalloc** — global allocator to reduce contention under high concurrency (1400+ tasks); page-agnostic, runs on aarch64 4K- and 64K-page kernels
+- **`Arc<str>` prompts** — zero-copy prompt sharing across tokio tasks, eliminating ~3 GB peak memory at 100k requests with 8k-token prompts
+- **Spawn-per-request** — `tokio::spawn` per request with a `Semaphore` for concurrency control (matches Python's asyncio pattern)
+- **rayon** — parallel dataset generation across CPU cores (200–500× faster than Python for 100k+ prompts)
+- **Enum dispatch** — backend variants instead of trait objects (avoids async trait-object limitations)
+- **Typed SSE deserialization** — `CompletionChunk`/`ChatChunk` structs skip unused JSON fields (cheaper than `serde_json::Value`)
+- **Speculative JSON parse** — SSE handler uses `serde_json::value::RawValue` to detect complete JSON before `\n\n` arrives, improving TTFT/ITL accuracy when TCP segments split
+- **Connection error retry** — automatic retry with backoff on connection reset/timeout/refused (up to 3 attempts)
+- **Tokenizer verification cache** — server-side token-length verification is cached per model+server pair
+
+### Behavioral parity with Python
+
+The Rust implementation matches Python `vllm bench serve` in:
+
+- SSE streaming protocol handling (including speculative parse for split TCP segments)
+- Timing semantics (monotonic `Instant` matching Python's `time.perf_counter()`)
+- Chat vs. completions differences (`max_completion_tokens` vs. `max_tokens`, Content-Type, timestamp placement)
+- JSON output schema (all fields, key naming, `request_rate` as the string `"inf"`)
+- Rate control (Gamma distribution, normalization, burstiness, linear/exponential ramp-up)
+- Metrics (TTFT/TPOT/ITL/E2EL percentiles, peak tokens/sec, peak concurrency, goodput)
+- Sampling parameters merged into the request body via `extra_body` (same precedence rules)
+
+## Environment Variables
+
+| Variable | Description |
+| ---------- | ------------- |
+| `OPENAI_API_KEY` | API key for authenticated endpoints (cached, not read per-request) |
+| `HF_TOKEN` | HuggingFace token for gated model tokenizers and gated datasets |
+| `TOKIO_WORKER_THREADS` | Override tokio worker thread count (default: physical cores) |
+
+## License
+
+Apache-2.0
+</content>
+</invoke>
@@ -0,0 +1,213 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+pub mod openai_chat;
+pub mod openai_completions;
+pub mod pooling;
+pub mod streaming;
+
+use std::collections::HashMap;
+use std::sync::Arc;
+
+use serde::{Deserialize, Serialize};
+
+// --- Typed SSE chunk structs for zero-alloc deserialization ---
+// Using typed deserialization avoids building a full serde_json::Value tree.
+// Only the fields we need are extracted; everything else is skipped by serde.
+
+/// Completions API streaming chunk (minimal fields).
+#[derive(Deserialize)]
+pub struct CompletionChunk {
+    #[serde(default)]
+    pub choices: Vec<CompletionChoice>,
+    pub usage: Option<ChunkUsage>,
+}
+
+#[derive(Deserialize)]
+pub struct CompletionChoice {
+    pub text: Option<String>,
+}
+
+/// Chat API streaming chunk (minimal fields).
+#[derive(Deserialize)]
+pub struct ChatChunk {
+    #[serde(default)]
+    pub choices: Vec<ChatChoice>,
+    pub usage: Option<ChunkUsage>,
+}
+
+#[derive(Deserialize)]
+pub struct ChatChoice {
+    pub delta: Option<ChatDelta>,
+}
+
+#[derive(Deserialize)]
+pub struct ChatDelta {
+    pub content: Option<String>,
+}
+
+#[derive(Deserialize)]
+pub struct ChunkUsage {
+    pub completion_tokens: Option<u64>,
+}
+
+use crate::cli::BackendKind;
+use crate::error::Result;
+
+/// Input for a single benchmark request.
+#[derive(Debug, Clone)]
+pub struct RequestFuncInput {
+    pub prompt: Arc<str>,
+    pub api_url: String,
+    pub prompt_len: usize,
+    pub output_len: usize,
+    pub model: String,
+    pub model_name: Option<String>,
+    pub logprobs: Option<usize>,
+    pub extra_headers: Option<HashMap<String, String>>,
+    pub extra_body: Option<serde_json::Value>,
+    pub ignore_eos: bool,
+    pub request_id: Option<String>,
+    /// Pre-built messages array for multi-turn conversations.
+    /// When set, the chat backend uses this instead of building from `prompt`.
+    pub messages: Option<serde_json::Value>,
+    /// Pre-computed token IDs for this prompt.
+    /// When set, the completions backend sends these directly via `prompt_token_ids`
+    /// instead of the text `prompt`, skipping server-side tokenization.
+    pub prompt_token_ids: Option<Arc<[u32]>>,
+    /// Multimodal content as pre-serialized JSON fragments.
+    /// When set, the chat backend concatenates these directly into the payload bytes,
+    /// avoiding any parsing or deep-cloning of base64 image data.
+    pub multi_modal_content: Option<Arc<[Arc<str>]>>,
+    /// Complete pre-serialized chat `messages` array (--enable-multimodal-chat).
+    /// When set, the chat backend splices it verbatim into the payload bytes,
+    /// taking precedence over `messages`, `prompt`, and `multi_modal_content`.
+    pub chat_messages_json: Option<Arc<str>>,
+    /// Multiple text inputs for one request (pooling backends only):
+    /// embeddings batch (`"input": [...]`) or rerank query+documents.
+    pub prompt_list: Option<Arc<[Arc<str>]>>,
+}
+
+/// Output from a single benchmark request including timing metrics.
+#[derive(Debug, Clone, Serialize, Deserialize)]
+pub struct RequestFuncOutput {
+    pub generated_text: String,
+    pub success: bool,
+    pub latency: f64,
+    pub output_tokens: usize,
+    pub ttft: f64,
+    pub itl: Vec<f64>,
+    pub tpot: f64,
+    pub prompt_len: usize,
+    pub error: String,
+    pub start_time: f64,
+}
+
+impl Default for RequestFuncOutput {
+    fn default() -> Self {
+        Self {
+            generated_text: String::new(),
+            success: false,
+            latency: 0.0,
+            output_tokens: 0,
+            ttft: 0.0,
+            itl: Vec::new(),
+            tpot: 0.0,
+            prompt_len: 0,
+            error: String::new(),
+            start_time: 0.0,
+        }
+    }
+}
+
+impl Default for RequestFuncInput {
+    fn default() -> Self {
+        Self {
+            prompt: Arc::from(""),
+            api_url: String::new(),
+            prompt_len: 0,
+            output_len: 0,
+            model: String::new(),
+            model_name: None,
+            logprobs: None,
+            extra_headers: None,
+            extra_body: None,
+            ignore_eos: false,
+            request_id: None,
+            messages: None,
+            prompt_token_ids: None,
+            multi_modal_content: None,
+            chat_messages_json: None,
+            prompt_list: None,
+        }
+    }
+}
+
+/// Enum dispatch for backend implementations (avoids async trait object issues).
+#[derive(Clone)]
+pub enum Backend {
+    OpenAICompletions(openai_completions::OpenAICompletionsBackend),
+    OpenAIChat(openai_chat::OpenAIChatBackend),
+    Pooling(pooling::PoolingBackend),
+}
+
+impl Backend {
+    /// Send a single request and collect timing metrics.
+    pub async fn send_request(
+        &self,
+        input: &RequestFuncInput,
+        client: &reqwest::Client,
+    ) -> Result<RequestFuncOutput> {
+        match self {
+            Backend::OpenAICompletions(b) => b.send_request(input, client).await,
+            Backend::OpenAIChat(b) => b.send_request(input, client).await,
+            Backend::Pooling(b) => b.send_request(input, client).await,
+        }
+    }
+}
+
+/// Get a backend by kind.
+pub fn get_backend(kind: BackendKind) -> Result<Backend> {
+    match kind {
+        BackendKind::Vllm | BackendKind::Openai => Ok(Backend::OpenAICompletions(
+            openai_completions::OpenAICompletionsBackend,
+        )),
+        BackendKind::OpenaiChat => Ok(Backend::OpenAIChat(openai_chat::OpenAIChatBackend)),
+        kind if kind.is_pooling() => Ok(Backend::Pooling(pooling::PoolingBackend { kind })),
+        _ => unreachable!(),
+    }
+}
+
+/// Cached API key to avoid per-request env var syscall.
+static API_KEY: std::sync::OnceLock<Option<String>> = std::sync::OnceLock::new();
+
+fn cached_api_key() -> &'static Option<String> {
+    API_KEY.get_or_init(|| std::env::var("OPENAI_API_KEY").ok())
+}
+
+/// Build common headers including auth and extras.
+pub fn build_headers(
+    content_type: Option<&str>,
+    extra_headers: &Option<HashMap<String, String>>,
+    request_id: &Option<String>,
+) -> HashMap<String, String> {
+    let mut headers = HashMap::new();
+
+    if let Some(ct) = content_type {
+        headers.insert("Content-Type".to_string(), ct.to_string());
+    }
+
+    if let Some(api_key) = cached_api_key() {
+        headers.insert("Authorization".to_string(), format!("Bearer {api_key}"));
+    }
+
+    if let Some(extra) = extra_headers {
+        headers.extend(extra.iter().map(|(k, v)| (k.clone(), v.clone())));
+    }
+
+    if let Some(rid) = request_id {
+        headers.insert("x-request-id".to_string(), rid.clone());
+    }
+
+    headers
+}
@@ -0,0 +1,376 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+use std::time::Instant;
+
+use futures::StreamExt;
+
+use super::streaming::{StreamedResponseHandler, trim_bytes};
+use super::{ChatChunk, RequestFuncInput, RequestFuncOutput, build_headers};
+use crate::error::Result;
+
+/// Backend for OpenAI Chat Completions API (/v1/chat/completions).
+#[derive(Clone)]
+pub struct OpenAIChatBackend;
+
+impl OpenAIChatBackend {
+    pub async fn send_request(
+        &self,
+        input: &RequestFuncInput,
+        client: &reqwest::Client,
+    ) -> Result<RequestFuncOutput> {
+        // Content-Type is set below by `.json()` / `.header()`; keep it out of
+        // headers_map to avoid a duplicate that strict gateways reject.
+        let headers_map = build_headers(None, &input.extra_headers, &input.request_id);
+
+        let mut output = RequestFuncOutput {
+            prompt_len: input.prompt_len,
+            itl: Vec::with_capacity(input.output_len.max(1)),
+            ..Default::default()
+        };
+
+        let st = Instant::now();
+
+        let mut most_recent_timestamp = st;
+        let mut generated_text = String::new();
+        let mut first_token_received = false;
+
+        // Build request: use zero-copy raw JSON for multimodal, serde_json for text-only
+        let mut request =
+            if input.multi_modal_content.is_some() || input.chat_messages_json.is_some() {
+                let payload_bytes = build_mm_payload(input);
+                client
+                    .post(&input.api_url)
+                    .header("content-type", "application/json")
+                    .body(payload_bytes)
+            } else {
+                let payload = build_text_payload(input);
+                client.post(&input.api_url).json(&payload)
+            };
+        for (k, v) in &headers_map {
+            request = request.header(k, v);
+        }
+
+        match request.send().await {
+            Ok(response) => {
+                if response.status().is_success() {
+                    let mut handler = StreamedResponseHandler::new();
+                    let mut stream = response.bytes_stream();
+
+                    while let Some(chunk_result) = stream.next().await {
+                        let chunk_bytes = match chunk_result {
+                            Ok(b) => b,
+                            Err(e) => {
+                                output.success = false;
+                                output.error = format!("Stream error: {e}");
+                                return Ok(output);
+                            }
+                        };
+
+                        let trimmed_bytes = trim_bytes(&chunk_bytes);
+                        if trimmed_bytes.is_empty() {
+                            continue;
+                        }
+
+                        let messages = handler.add_chunk(trimmed_bytes);
+                        for message in messages {
+                            // Skip SSE comments
+                            if message.starts_with(':') {
+                                continue;
+                            }
+
+                            // Handle multi-field SSE events (e.g., Dynamo sends
+                            // "event: message\ndata: {...}"). Extract the data: line.
+                            let raw = if message.contains('\n') {
+                                match message.lines().find(|l| l.starts_with("data: ")) {
+                                    Some(l) => l,
+                                    None => continue,
+                                }
+                            } else {
+                                message.as_str()
+                            };
+
+                            let chunk = raw.strip_prefix("data: ").unwrap_or(raw);
+
+                            if chunk == "[DONE]" {
+                                continue;
+                            }
+
+                            // Python chat backend: timestamp is captured for ALL
+                            // non-DONE messages, and most_recent_timestamp is updated
+                            // unconditionally (outside `if choices:`). This differs from
+                            // completions which only timestamps content chunks.
+                            let timestamp = Instant::now();
+
+                            let data: ChatChunk = match serde_json::from_str(chunk) {
+                                Ok(d) => d,
+                                Err(_) => continue,
+                            };
+
+                            if !data.choices.is_empty() {
+                                let content = data.choices[0]
+                                    .delta
+                                    .as_ref()
+                                    .and_then(|d| d.content.as_deref())
+                                    .unwrap_or("");
+
+                                if !first_token_received {
+                                    first_token_received = true;
+                                    output.ttft = timestamp.duration_since(st).as_secs_f64();
+                                } else {
+                                    output.itl.push(
+                                        timestamp
+                                            .duration_since(most_recent_timestamp)
+                                            .as_secs_f64(),
+                                    );
+                                }
+
+                                generated_text.push_str(content);
+                            }
+                            // Separate `if` (not `else if`) — Dynamo may send
+                            // both choices and usage in the same chunk.
+                            if let Some(ref usage) = data.usage
+                                && let Some(ct) = usage.completion_tokens
+                            {
+                                output.output_tokens = ct as usize;
+                            }
+
+                            most_recent_timestamp = timestamp;
+                        }
+                    }
+
+                    output.generated_text = generated_text;
+                    output.success = true;
+                    output.latency = most_recent_timestamp.duration_since(st).as_secs_f64();
+                } else {
+                    let status = response.status();
+                    let body = response.text().await.unwrap_or_default();
+                    output.error = if body.is_empty() {
+                        format!("HTTP {status}")
+                    } else {
+                        format!("HTTP {status}: {body}")
+                    };
+                    output.success = false;
+                }
+            }
+            Err(e) => {
+                output.success = false;
+                output.error = format!("{e:#}");
+            }
+        }
+
+        Ok(output)
+    }
+}
+
+/// Build a JSON payload for text-only (non-multimodal) requests using serde_json.
+fn build_text_payload(input: &RequestFuncInput) -> serde_json::Value {
+    let model = input.model_name.as_deref().unwrap_or(&input.model);
+
+    let messages = if let Some(ref msgs) = input.messages {
+        msgs.clone()
+    } else {
+        let content = serde_json::json!([
+            {"type": "text", "text": input.prompt}
+        ]);
+        serde_json::json!([{"role": "user", "content": content}])
+    };
+
+    let mut payload = serde_json::json!({
+        "model": model,
+        "messages": messages,
+        "max_completion_tokens": input.output_len,
+        "stream": true,
+        "stream_options": {
+            "include_usage": true,
+        },
+    });
+
+    if input.ignore_eos {
+        payload["ignore_eos"] = serde_json::json!(true);
+    }
+    if let Some(serde_json::Value::Object(map)) = input.extra_body.as_ref() {
+        for (k, v) in map {
+            payload[k] = v.clone();
+        }
+    }
+
+    payload
+}
+
+/// Build the JSON payload as raw bytes for multimodal requests.
+///
+/// This is the zero-copy fast path: pre-serialized mm content fragments
+/// (each ~200KB+ of base64 image data) are concatenated directly into the
+/// output buffer without being parsed, cloned, or re-serialized.
+///
+/// Saves ~200KB of allocation + copy per image per request compared to
+/// the serde_json::Value approach.
+fn build_mm_payload(input: &RequestFuncInput) -> Vec<u8> {
+    let model = input.model_name.as_deref().unwrap_or(&input.model);
+
+    // Estimate total size: JSON overhead (~300 bytes) + prompt + mm fragments
+    let mm_total: usize = input
+        .multi_modal_content
+        .as_ref()
+        .map(|mm| mm.iter().map(|f| f.len() + 1).sum())
+        .unwrap_or(0)
+        + input.chat_messages_json.as_ref().map_or(0, |m| m.len());
+    let estimated = 512 + input.prompt.len() * 2 + mm_total;
+    let mut json = String::with_capacity(estimated);
+
+    // {"model": <model>
+    json.push_str(r#"{"model":"#);
+    // serde_json::to_string on &str produces a JSON-escaped quoted string
+    json.push_str(&serde_json::to_string(model).unwrap());
+
+    json.push_str(r#","messages":"#);
+    if let Some(ref msgs) = input.chat_messages_json {
+        // --enable-multimodal-chat: the dataset pre-built the full messages
+        // array (text + mm parts); splice it verbatim.
+        json.push_str(msgs);
+    } else {
+        let mm = input.multi_modal_content.as_ref().unwrap();
+
+        // [{"role":"user","content":[ <text part>
+        json.push_str(r#"[{"role":"user","content":[{"type":"text","text":""#);
+        // JSON-escape the prompt text (handles \n, \t, unicode, quotes)
+        push_json_escaped_str(&mut json, &input.prompt);
+        json.push_str(r#""}"#);
+
+        // ,<mm fragment 1>,<mm fragment 2>,...
+        for fragment in mm.iter() {
+            json.push(',');
+            json.push_str(fragment);
+        }
+
+        // Close content, message, messages
+        json.push_str(r#"]}]"#);
+    }
+
+    // ,"max_completion_tokens": N, "stream": true, ...
+    json.push_str(r##","max_completion_tokens":"##);
+    json.push_str(&input.output_len.to_string());
+    json.push_str(r##","stream":true,"stream_options":{"include_usage":true}"##);
+
+    if input.ignore_eos {
+        json.push_str(r#","ignore_eos":true"#);
+    }
+
+    // Merge extra_body key-value pairs, skipping keys already set above
+    if let Some(serde_json::Value::Object(map)) = input.extra_body.as_ref() {
+        for (k, v) in map {
+            match k.as_str() {
+                "model"
+                | "messages"
+                | "max_completion_tokens"
+                | "stream"
+                | "stream_options"
+                | "ignore_eos" => continue,
+                _ => {
+                    json.push(',');
+                    json.push_str(&serde_json::to_string(k).unwrap());
+                    json.push(':');
+                    json.push_str(&serde_json::to_string(v).unwrap());
+                }
+            }
+        }
+    }
+
+    json.push('}');
+    json.into_bytes()
+}
+
+/// Write a JSON-escaped string (without surrounding quotes) into the buffer.
+///
+/// Handles: `\n`, `\r`, `\t`, `\\`, `\"`, and control characters.
+/// This avoids the allocation of `serde_json::to_string` which produces
+/// a new String with surrounding quotes.
+fn push_json_escaped_str(buf: &mut String, s: &str) {
+    use std::fmt::Write;
+    for ch in s.chars() {
+        match ch {
+            '"' => buf.push_str(r#"\""#),
+            '\\' => buf.push_str(r"\\"),
+            '\n' => buf.push_str(r"\n"),
+            '\r' => buf.push_str(r"\r"),
+            '\t' => buf.push_str(r"\t"),
+            c if c.is_control() => {
+                // \uXXXX escape for control characters
+                for unit in c.encode_utf16(&mut [0; 2]) {
+                    write!(buf, "\\u{unit:04x}").unwrap();
+                }
+            }
+            c => buf.push(c),
+        }
+    }
+}
+
+#[cfg(test)]
+mod tests {
+    use std::sync::Arc;
+
+    use super::*;
+
+    fn mm_input() -> RequestFuncInput {
+        let frag: Arc<str> =
+            Arc::from(r#"{"type":"image_url","image_url":{"url":"data:image/jpeg;base64,AAAA"}}"#);
+        RequestFuncInput {
+            prompt: Arc::from("hello \"world\"\nline2"),
+            model: "test-model".to_string(),
+            output_len: 128,
+            multi_modal_content: Some(Arc::from(vec![frag])),
+            ..Default::default()
+        }
+    }
+
+    /// Regression test: the assembled multimodal payload must be valid JSON
+    /// (the text part once shipped without its opening quote — a raw-string
+    /// delimiter eating the trailing `"` in `"text":"`).
+    #[test]
+    fn test_build_mm_payload_is_valid_json() {
+        let payload = build_mm_payload(&mm_input());
+        let v: serde_json::Value =
+            serde_json::from_slice(&payload).expect("mm payload must be valid JSON");
+        assert_eq!(v["model"], "test-model");
+        assert_eq!(v["messages"][0]["role"], "user");
+        let content = v["messages"][0]["content"].as_array().unwrap();
+        assert_eq!(content[0]["type"], "text");
+        assert_eq!(content[0]["text"], "hello \"world\"\nline2");
+        assert_eq!(content[1]["type"], "image_url");
+        assert_eq!(v["max_completion_tokens"], 128);
+        assert_eq!(v["stream"], true);
+        assert_eq!(v["stream_options"]["include_usage"], true);
+    }
+
+    /// --enable-multimodal-chat (dataset pre-built messages) must produce a
+    /// payload semantically identical to the fragment-assembly path.
+    #[test]
+    fn test_chat_messages_json_path_equivalent_to_fragment_path() {
+        let base = mm_input();
+        let fragment_payload = build_mm_payload(&base);
+
+        let mut chat = base.clone();
+        let mm = chat.multi_modal_content.take().unwrap();
+        let msgs = crate::datasets::random_mm::build_chat_messages_json(&chat.prompt, Some(&mm));
+        chat.chat_messages_json = Some(Arc::from(msgs.as_str()));
+        let chat_payload = build_mm_payload(&chat);
+
+        let a: serde_json::Value = serde_json::from_slice(&fragment_payload).unwrap();
+        let b: serde_json::Value = serde_json::from_slice(&chat_payload).unwrap();
+        assert_eq!(a, b);
+    }
+
+    /// ignore_eos and extra_body must survive the raw-splice path.
+    #[test]
+    fn test_mm_payload_tail_fields() {
+        let mut input = mm_input();
+        input.ignore_eos = true;
+        input.extra_body = Some(serde_json::json!({"temperature": 0.5, "stream": false}));
+        let v: serde_json::Value = serde_json::from_slice(&build_mm_payload(&input)).unwrap();
+        assert_eq!(v["ignore_eos"], true);
+        assert_eq!(v["temperature"], 0.5);
+        // keys already set above must not be overridden by extra_body
+        assert_eq!(v["stream"], true);
+    }
+}
@@ -0,0 +1,199 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+use std::error::Error as StdError;
+use std::time::Instant;
+
+use futures::StreamExt;
+
+use super::streaming::{StreamedResponseHandler, trim_bytes};
+use super::{CompletionChunk, RequestFuncInput, RequestFuncOutput, build_headers};
+use crate::error::Result;
+
+/// Backend for OpenAI-compatible Completions API (/v1/completions).
+/// Used by "vllm" and "openai" backends.
+#[derive(Clone)]
+pub struct OpenAICompletionsBackend;
+
+impl OpenAICompletionsBackend {
+    pub async fn send_request(
+        &self,
+        input: &RequestFuncInput,
+        client: &reqwest::Client,
+    ) -> Result<RequestFuncOutput> {
+        let model = input.model_name.as_deref().unwrap_or(&input.model);
+
+        // When prompt_token_ids are available, send them as the `prompt` value
+        // (JSON array of integers). vLLM's completions API accepts both string
+        // and token ID array as `prompt`, skipping server-side tokenization.
+        let prompt_value = if let Some(ref token_ids) = input.prompt_token_ids {
+            serde_json::json!(token_ids.as_ref())
+        } else {
+            serde_json::json!(input.prompt)
+        };
+
+        let mut payload = serde_json::json!({
+            "model": model,
+            "prompt": prompt_value,
+            "max_tokens": input.output_len,
+            "stream": true,
+            "stream_options": {
+                "include_usage": true,
+            },
+        });
+
+        // Always include logprobs (null when not set) — matches Python which
+        // sends logprobs=None explicitly rather than omitting the key.
+        payload["logprobs"] = match input.logprobs {
+            Some(n) => serde_json::json!(n),
+            None => serde_json::Value::Null,
+        };
+
+        // Apply ignore_eos and extra_body
+        if input.ignore_eos {
+            payload["ignore_eos"] = serde_json::json!(true);
+        }
+        if let Some(serde_json::Value::Object(map)) = input.extra_body.as_ref() {
+            for (k, v) in map {
+                payload[k] = v.clone();
+            }
+        }
+
+        let headers_map = build_headers(None, &input.extra_headers, &input.request_id);
+
+        let mut output = RequestFuncOutput {
+            prompt_len: input.prompt_len,
+            itl: Vec::with_capacity(input.output_len.max(1)),
+            ..Default::default()
+        };
+
+        let st = Instant::now();
+        // start_time is overwritten by benchmark.rs with monotonic offset
+
+        let mut most_recent_timestamp = st;
+        let mut generated_text = String::new();
+        let mut first_chunk_received = false;
+
+        let mut request = client.post(&input.api_url).json(&payload);
+        for (k, v) in &headers_map {
+            request = request.header(k, v);
+        }
+
+        match request.send().await {
+            Ok(response) => {
+                if response.status().is_success() {
+                    let mut handler = StreamedResponseHandler::new();
+                    let mut stream = response.bytes_stream();
+
+                    while let Some(chunk_result) = stream.next().await {
+                        let chunk_bytes = match chunk_result {
+                            Ok(b) => b,
+                            Err(e) => {
+                                output.success = false;
+                                output.error = format!("Stream error: {e}");
+                                return Ok(output);
+                            }
+                        };
+
+                        let trimmed_bytes = trim_bytes(&chunk_bytes);
+                        if trimmed_bytes.is_empty() {
+                            continue;
+                        }
+
+                        let messages = handler.add_chunk(trimmed_bytes);
+                        for message in messages {
+                            // Skip SSE comments
+                            if message.starts_with(':') {
+                                continue;
+                            }
+
+                            // Handle multi-field SSE events (e.g., Dynamo sends
+                            // "event: message\ndata: {...}"). Extract the data: line.
+                            let raw = if message.contains('\n') {
+                                match message.lines().find(|l| l.starts_with("data: ")) {
+                                    Some(l) => l,
+                                    None => continue,
+                                }
+                            } else {
+                                message.as_str()
+                            };
+
+                            let chunk = raw.strip_prefix("data: ").unwrap_or(raw);
+
+                            if chunk == "[DONE]" {
+                                continue;
+                            }
+
+                            // Typed deserialization — avoids allocating a full
+                            // serde_json::Value tree; only extracts needed fields.
+                            let data: CompletionChunk = match serde_json::from_str(chunk) {
+                                Ok(d) => d,
+                                Err(_) => continue,
+                            };
+
+                            if !data.choices.is_empty() {
+                                let text = data.choices[0].text.as_deref().unwrap_or("");
+
+                                let timestamp = Instant::now();
+
+                                if !first_chunk_received {
+                                    first_chunk_received = true;
+                                    output.ttft = timestamp.duration_since(st).as_secs_f64();
+                                } else {
+                                    output.itl.push(
+                                        timestamp
+                                            .duration_since(most_recent_timestamp)
+                                            .as_secs_f64(),
+                                    );
+                                }
+
+                                most_recent_timestamp = timestamp;
+                                generated_text.push_str(text);
+                            }
+                            // Separate `if` (not `else if`) — Dynamo may send
+                            // both choices and usage in the same chunk.
+                            if let Some(ref usage) = data.usage
+                                && let Some(ct) = usage.completion_tokens
+                            {
+                                output.output_tokens = ct as usize;
+                            }
+                        }
+                    }
+
+                    if first_chunk_received {
+                        output.success = true;
+                    } else {
+                        output.success = false;
+                        output.error = "Never received a valid chunk to calculate TTFT. \
+                                        This response will be marked as failed!"
+                            .to_string();
+                    }
+                    output.generated_text = generated_text;
+                    output.latency = most_recent_timestamp.duration_since(st).as_secs_f64();
+                } else {
+                    let status = response.status();
+                    let body = response.text().await.unwrap_or_default();
+                    output.error = if body.is_empty() {
+                        format!("HTTP {status}")
+                    } else {
+                        format!("HTTP {status}: {body}")
+                    };
+                    output.success = false;
+                }
+            }
+            Err(e) => {
+                output.success = false;
+                // Capture full error chain for debugging
+                let mut error_msg = format!("{e}");
+                let mut source = e.source();
+                while let Some(cause) = source {
+                    error_msg.push_str(&format!("\n  Caused by: {cause}"));
+                    source = cause.source();
+                }
+                output.error = error_msg;
+            }
+        }
+
+        Ok(output)
+    }
+}
@@ -0,0 +1,325 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+//! Pooling/embedding backends: non-streaming HTTP POST for embedding, pooling, and rerank
+//! endpoints.
+//!
+//! Supported variants:
+//! - `openai-embeddings`: Standard OpenAI `/v1/embeddings` with text input
+//! - `openai-embeddings-chat`: OpenAI `/v1/embeddings` with chat message format (supports
+//!   multimodal)
+//! - `vllm-pooling`: vLLM `/v1/pooling` endpoint
+//! - `vllm-rerank`: vLLM `/v1/rerank` endpoint (query + documents)
+
+use std::time::Instant;
+
+use crate::backends::{RequestFuncInput, RequestFuncOutput, build_headers};
+use crate::cli::BackendKind;
+use crate::error::Result;
+
+/// Response from embedding/pooling endpoints (minimal fields for usage extraction).
+#[derive(serde::Deserialize)]
+struct PoolingResponse {
+    usage: Option<PoolingUsage>,
+}
+
+#[derive(serde::Deserialize)]
+struct PoolingUsage {
+    prompt_tokens: Option<u64>,
+}
+
+#[derive(Clone)]
+pub struct PoolingBackend {
+    pub kind: BackendKind,
+}
+
+impl PoolingBackend {
+    pub async fn send_request(
+        &self,
+        input: &RequestFuncInput,
+        client: &reqwest::Client,
+    ) -> Result<RequestFuncOutput> {
+        // Preserve client-side prompt_len as fallback if server doesn't report usage.
+        let mut output = RequestFuncOutput {
+            prompt_len: input.prompt_len,
+            ..Default::default()
+        };
+
+        let headers = build_headers(
+            Some("application/json"),
+            &input.extra_headers,
+            &input.request_id,
+        );
+
+        let payload = self.build_payload(input);
+
+        let mut request = client.post(&input.api_url);
+        for (k, v) in &headers {
+            request = request.header(k, v);
+        }
+
+        let st = Instant::now();
+
+        let response = match request.json(&payload).send().await {
+            Ok(r) => r,
+            Err(e) => {
+                output.error = format!("Request failed: {e}");
+                return Ok(output);
+            }
+        };
+
+        if response.status().is_success() {
+            let latency = st.elapsed().as_secs_f64();
+            output.latency = latency;
+            output.ttft = latency;
+            output.success = true;
+
+            // Parse usage from response; keep client-side prompt_len as fallback.
+            match response.json::<PoolingResponse>().await {
+                Ok(data) => {
+                    if let Some(usage) = data.usage
+                        && let Some(tokens) = usage.prompt_tokens
+                    {
+                        output.prompt_len = tokens as usize;
+                    }
+                }
+                Err(_) => {
+                    // Response parsed but no usage — keep client-side prompt_len
+                }
+            }
+        } else {
+            let status = response.status();
+            let body = response.text().await.unwrap_or_default();
+            output.error = format!("HTTP {status}: {body}");
+        }
+
+        Ok(output)
+    }
+
+    fn build_payload(&self, input: &RequestFuncInput) -> serde_json::Value {
+        let model = input.model_name.as_deref().unwrap_or(&input.model);
+
+        // For "input" field (openai-embeddings, vllm-pooling): a batched request
+        // (--random-batch-size) sends the text list; otherwise prefer prompt_token_ids
+        // when available. The random dataset sets prompt="" and relies on token IDs;
+        // the OpenAI embeddings API accepts both text strings and token ID arrays.
+        // Note: embeddings-chat uses text in messages; vllm-rerank uses text as query.
+        let input_value = if let Some(ref list) = input.prompt_list {
+            serde_json::json!(list.iter().map(|s| s.as_ref()).collect::<Vec<&str>>())
+        } else if let Some(ref token_ids) = input.prompt_token_ids {
+            serde_json::json!(token_ids.as_ref())
+        } else {
+            serde_json::json!(input.prompt.as_ref())
+        };
+
+        let is_vllm_backend = matches!(
+            self.kind,
+            BackendKind::VllmPooling | BackendKind::VllmRerank
+        );
+
+        let mut payload = match self.kind {
+            BackendKind::OpenaiEmbeddings => {
+                let mut p = serde_json::json!({
+                    "model": model,
+                    "input": input_value,
+                });
+                // truncate_prompt_tokens is vLLM-specific; only include for vLLM backends
+                // to avoid breaking standard OpenAI providers.
+                if is_vllm_backend {
+                    p["truncate_prompt_tokens"] = serde_json::json!(-1);
+                }
+                p
+            }
+            BackendKind::OpenaiEmbeddingsChat => {
+                // Chat format: uses text prompt in messages array (for multimodal support).
+                // Python's _get_chat_content always returns a content array.
+                // Use raw string concatenation for multimodal fragments (zero-copy,
+                // avoids re-parsing ~200KB+ base64 per image).
+                let content_json = build_chat_content_json(input);
+
+                let mut p = serde_json::json!({
+                    "model": model,
+                    "messages": [{"role": "user", "content": content_json}],
+                });
+                if is_vllm_backend {
+                    p["truncate_prompt_tokens"] = serde_json::json!(-1);
+                }
+                p
+            }
+            BackendKind::VllmPooling => {
+                serde_json::json!({
+                    "model": model,
+                    "input": input_value,
+                    "truncate_prompt_tokens": -1,
+                })
+            }
+            BackendKind::VllmRerank => {
+                // random-rerank dataset: prompt_list = [query, doc1, doc2, ...]
+                // (mirrors Python async_request_vllm_rerank).
+                if let Some(ref list) = input.prompt_list {
+                    if list.len() < 2 {
+                        eprintln!(
+                            "WARNING: vllm-rerank request has no documents \
+                             (prompt_list needs [query, doc, ...])"
+                        );
+                    }
+                    let query = list.first().map(|s| s.as_ref()).unwrap_or("");
+                    let documents: Vec<&str> = list.iter().skip(1).map(|s| s.as_ref()).collect();
+                    serde_json::json!({
+                        "model": model,
+                        "query": query,
+                        "documents": documents,
+                        "truncate_prompt_tokens": -1,
+                    })
+                } else {
+                    // Legacy path: text prompt as query, documents via --extra-body.
+                    let query = input.prompt.as_ref();
+                    if query.is_empty() && input.prompt_token_ids.is_some() {
+                        eprintln!(
+                            "WARNING: vllm-rerank received empty query (random dataset uses \
+                             token IDs only). Use --dataset-name random-rerank for meaningful \
+                             rerank benchmarks."
+                        );
+                    }
+                    serde_json::json!({
+                        "model": model,
+                        "query": query,
+                        "truncate_prompt_tokens": -1,
+                    })
+                }
+            }
+            _ => unreachable!("PoolingBackend with non-pooling kind"),
+        };
+
+        // Merge extra_body fields into payload
+        if let Some(ref extra) = input.extra_body
+            && let (Some(base), Some(extra_obj)) = (payload.as_object_mut(), extra.as_object())
+        {
+            for (k, v) in extra_obj {
+                base.insert(k.clone(), v.clone());
+            }
+        }
+
+        payload
+    }
+}
+
+/// Build the chat content JSON array for embeddings-chat.
+/// Uses raw string concatenation for multimodal fragments to avoid
+/// re-parsing large base64 image data (matching openai_chat.rs approach).
+fn build_chat_content_json(input: &RequestFuncInput) -> serde_json::Value {
+    if input.multi_modal_content.is_none() {
+        // Text-only: return content array with single text element
+        return serde_json::json!([{
+            "type": "text",
+            "text": input.prompt.as_ref(),
+        }]);
+    }
+
+    // Multimodal: build JSON string manually for zero-copy fragment embedding
+    let mm = input.multi_modal_content.as_ref().unwrap();
+    let prompt = input.prompt.as_ref();
+
+    let mm_total: usize = mm.iter().map(|f| f.len() + 1).sum();
+    let mut json = String::with_capacity(64 + prompt.len() * 2 + mm_total);
+
+    // [{"type":"text","text":"<prompt>"}
+    json.push_str(r#"[{"type":"text","text":""#);
+    push_json_escaped_str(&mut json, prompt);
+    json.push_str(r#""}"#);
+
+    // ,<mm fragment 1>,<mm fragment 2>,...
+    for fragment in mm.iter() {
+        json.push(',');
+        json.push_str(fragment);
+    }
+
+    json.push(']');
+
+    // Parse the assembled string into a Value for embedding in the payload.
+    // This parse is O(n) but operates on the pre-built string once, not per-fragment.
+    serde_json::from_str(&json).unwrap_or_else(|_| {
+        serde_json::json!([{
+            "type": "text",
+            "text": input.prompt.as_ref(),
+        }])
+    })
+}
+
+/// Escape a string for safe JSON embedding (matching openai_chat.rs).
+fn push_json_escaped_str(buf: &mut String, s: &str) {
+    use std::fmt::Write;
+    for ch in s.chars() {
+        match ch {
+            '"' => buf.push_str(r#"\""#),
+            '\\' => buf.push_str(r"\\"),
+            '\n' => buf.push_str(r"\n"),
+            '\r' => buf.push_str(r"\r"),
+            '\t' => buf.push_str(r"\t"),
+            c if c < '\x20' => {
+                let _ = write!(buf, "\\u{:04x}", c as u32);
+            }
+            c => buf.push(c),
+        }
+    }
+}
+
+#[cfg(test)]
+mod tests {
+    use std::sync::Arc;
+
+    use super::*;
+
+    fn list(items: &[&str]) -> Option<Arc<[Arc<str>]>> {
+        Some(items.iter().map(|s| Arc::from(*s)).collect())
+    }
+
+    #[test]
+    fn test_embeddings_payload_batched_input() {
+        let backend = PoolingBackend {
+            kind: BackendKind::OpenaiEmbeddings,
+        };
+        let input = RequestFuncInput {
+            model: "bge".to_string(),
+            prompt_list: list(&["t1", "t2", "t3"]),
+            ..Default::default()
+        };
+        let payload = backend.build_payload(&input);
+        assert_eq!(payload["input"], serde_json::json!(["t1", "t2", "t3"]));
+        assert_eq!(payload["model"], "bge");
+        // truncate_prompt_tokens is vLLM-specific and deliberately omitted for
+        // the plain OpenAI embeddings backend.
+        assert!(payload.get("truncate_prompt_tokens").is_none());
+    }
+
+    #[test]
+    fn test_rerank_payload_query_and_documents() {
+        let backend = PoolingBackend {
+            kind: BackendKind::VllmRerank,
+        };
+        let input = RequestFuncInput {
+            model: "reranker".to_string(),
+            prompt_list: list(&["the query", "doc a", "doc b"]),
+            ..Default::default()
+        };
+        let payload = backend.build_payload(&input);
+        assert_eq!(payload["query"], "the query");
+        assert_eq!(payload["documents"], serde_json::json!(["doc a", "doc b"]));
+        assert_eq!(payload["truncate_prompt_tokens"], -1);
+    }
+
+    #[test]
+    fn test_rerank_payload_legacy_single_prompt() {
+        let backend = PoolingBackend {
+            kind: BackendKind::VllmRerank,
+        };
+        let input = RequestFuncInput {
+            model: "reranker".to_string(),
+            prompt: Arc::from("query text"),
+            ..Default::default()
+        };
+        let payload = backend.build_payload(&input);
+        assert_eq!(payload["query"], "query text");
+        assert!(payload.get("documents").is_none());
+    }
+}
@@ -0,0 +1,151 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+/// SSE streaming response handler.
+///
+/// Accumulates incoming byte chunks and extracts complete SSE messages.
+/// Mirrors Python's `StreamedResponseHandler` from endpoint_request_func.py:22-60.
+pub struct StreamedResponseHandler {
+    buffer: String,
+    /// Reusable message buffer — avoids allocating a new Vec per `add_chunk` call.
+    messages: Vec<String>,
+}
+
+impl StreamedResponseHandler {
+    pub fn new() -> Self {
+        Self {
+            buffer: String::with_capacity(4096),
+            messages: Vec::with_capacity(4),
+        }
+    }
+
+    /// Add a chunk of bytes and return any complete SSE messages.
+    ///
+    /// The returned slice borrows from the handler and is valid until the next
+    /// `add_chunk` call.
+    pub fn add_chunk(&mut self, chunk_bytes: &[u8]) -> &[String] {
+        self.messages.clear();
+
+        let chunk_str = String::from_utf8_lossy(chunk_bytes);
+        self.buffer.push_str(&chunk_str);
+
+        // Split by double newlines (SSE message separator)
+        while let Some(pos) = self.buffer.find("\n\n") {
+            let message = self.buffer[..pos].trim().to_string();
+            // Efficiently remove consumed bytes by shifting remaining data
+            self.buffer.drain(..pos + 2);
+            if !message.is_empty() {
+                self.messages.push(message);
+            }
+        }
+
+        // Handle buffered data without trailing `\n\n`.
+        // Matches Python's speculative json.loads() in StreamedResponseHandler.
+        // This matters for TTFT/ITL accuracy: when a data message and its `\n\n`
+        // arrive in separate TCP segments, we want to emit the message at the
+        // first segment's arrival time, not the second.
+        //
+        // Also handles multi-field SSE events where the buffer may start with
+        // "event: ...\ndata: ..." (Dynamo frontend).
+        let data_start = if self.buffer.starts_with("data: ") {
+            Some(0)
+        } else {
+            // Look for a "data: " line in multi-field events
+            self.buffer.find("\ndata: ").map(|p| p + 1)
+        };
+        if let Some(offset) = data_start {
+            let content = self.buffer[offset + 6..].trim();
+            if content == "[DONE]"
+                || (!content.is_empty()
+                    && serde_json::from_str::<&serde_json::value::RawValue>(content).is_ok())
+            {
+                self.messages.push(self.buffer.trim().to_string());
+                self.buffer.clear();
+            }
+        }
+
+        &self.messages
+    }
+}
+
+/// Trim leading/trailing ASCII whitespace from a byte slice.
+pub fn trim_bytes(bytes: &[u8]) -> &[u8] {
+    let start = bytes.iter().position(|b| !b.is_ascii_whitespace()).unwrap_or(bytes.len());
+    let end = bytes
+        .iter()
+        .rposition(|b| !b.is_ascii_whitespace())
+        .map(|p| p + 1)
+        .unwrap_or(start);
+    &bytes[start..end]
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+
+    #[test]
+    fn test_basic_sse() {
+        let mut handler = StreamedResponseHandler::new();
+        let msgs =
+            handler.add_chunk(b"data: {\"choices\":[{\"text\":\"hi\"}]}\n\ndata: [DONE]\n\n");
+        assert_eq!(msgs.len(), 2);
+        assert!(msgs[0].contains("choices"));
+        assert!(msgs[1].contains("[DONE]"));
+    }
+
+    #[test]
+    fn test_split_chunks() {
+        let mut handler = StreamedResponseHandler::new();
+        let msgs1 = handler.add_chunk(b"data: {\"cho");
+        assert!(msgs1.is_empty());
+        let msgs2 = handler.add_chunk(b"ices\":[{\"text\":\"a\"}]}\n\n");
+        assert_eq!(msgs2.len(), 1);
+    }
+
+    #[test]
+    fn test_comment_lines() {
+        let mut handler = StreamedResponseHandler::new();
+        let msgs = handler.add_chunk(b": ping\n\ndata: {\"test\":1}\n\n");
+        assert_eq!(msgs.len(), 2);
+        assert!(msgs[0].starts_with(":"));
+    }
+
+    #[test]
+    fn test_done_without_newlines() {
+        let mut handler = StreamedResponseHandler::new();
+        let msgs = handler.add_chunk(b"data: [DONE]");
+        assert_eq!(msgs.len(), 1);
+        assert!(msgs[0].contains("[DONE]"));
+    }
+
+    #[test]
+    fn test_incomplete_json_in_buffer() {
+        let mut handler = StreamedResponseHandler::new();
+        let msgs = handler.add_chunk(b"data: {\"partial\":");
+        assert!(msgs.is_empty());
+        // Complete JSON without \n\n — speculative parse emits it
+        let msgs2 = handler.add_chunk(b"true}");
+        assert_eq!(msgs2.len(), 1);
+        assert!(msgs2[0].contains("partial"));
+    }
+
+    #[test]
+    fn test_multi_field_sse_event() {
+        // Dynamo frontend sends "event: message\ndata: {...}\n\n"
+        let mut handler = StreamedResponseHandler::new();
+        let msgs =
+            handler.add_chunk(b"event: message\ndata: {\"choices\":[{\"text\":\"hi\"}]}\n\n");
+        assert_eq!(msgs.len(), 1);
+        assert!(msgs[0].contains("choices"));
+        assert!(msgs[0].contains("event: message"));
+    }
+
+    #[test]
+    fn test_multi_field_sse_speculative_parse() {
+        // Multi-field event without trailing \n\n — speculative parse should emit it
+        let mut handler = StreamedResponseHandler::new();
+        let msgs = handler.add_chunk(b"event: message\ndata: {\"choices\":[{\"text\":\"hi\"}]}");
+        assert_eq!(msgs.len(), 1);
+        assert!(msgs[0].contains("choices"));
+    }
+}
@@ -0,0 +1,743 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+use std::fmt;
+
+/// Backend type for the benchmark endpoint.
+#[derive(clap::ValueEnum, Debug, Clone, Copy, PartialEq, Eq)]
+pub enum BackendKind {
+    #[value(name = "vllm")]
+    Vllm,
+    #[value(name = "openai")]
+    Openai,
+    #[value(name = "openai-chat")]
+    OpenaiChat,
+    #[value(name = "openai-embeddings")]
+    OpenaiEmbeddings,
+    #[value(name = "openai-embeddings-chat")]
+    OpenaiEmbeddingsChat,
+    #[value(name = "vllm-pooling")]
+    VllmPooling,
+    #[value(name = "vllm-rerank")]
+    VllmRerank,
+}
+
+impl BackendKind {
+    pub fn as_str(self) -> &'static str {
+        match self {
+            Self::Vllm => "vllm",
+            Self::Openai => "openai",
+            Self::OpenaiChat => "openai-chat",
+            Self::OpenaiEmbeddings => "openai-embeddings",
+            Self::OpenaiEmbeddingsChat => "openai-embeddings-chat",
+            Self::VllmPooling => "vllm-pooling",
+            Self::VllmRerank => "vllm-rerank",
+        }
+    }
+
+    /// Return true if the backend is compatible with OpenAI-style API and sampling parameters.
+    pub fn is_openai_compatible(self) -> bool {
+        match self {
+            Self::Vllm | Self::Openai | Self::OpenaiChat => true,
+            Self::OpenaiEmbeddings
+            | Self::OpenaiEmbeddingsChat
+            | Self::VllmPooling
+            | Self::VllmRerank => false,
+        }
+    }
+
+    /// Return true if the backend is a pooling/embedding backend (non-generative).
+    pub fn is_pooling(self) -> bool {
+        matches!(
+            self,
+            Self::OpenaiEmbeddings
+                | Self::OpenaiEmbeddingsChat
+                | Self::VllmPooling
+                | Self::VllmRerank
+        )
+    }
+}
+
+impl fmt::Display for BackendKind {
+    fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
+        f.write_str(self.as_str())
+    }
+}
+
+/// Dataset to benchmark with.
+#[derive(clap::ValueEnum, Debug, Clone, Copy, PartialEq, Eq)]
+pub enum DatasetName {
+    #[value(name = "random")]
+    Random,
+    #[value(name = "random-mm")]
+    RandomMm,
+    #[value(name = "sharegpt")]
+    ShareGpt,
+    #[value(name = "sonnet")]
+    Sonnet,
+    #[value(name = "speed-bench", alias = "speed_bench")]
+    SpeedBench,
+    #[value(name = "hf")]
+    Hf,
+    #[value(name = "custom")]
+    Custom,
+    #[value(name = "prefix_repetition", alias = "prefix-repetition")]
+    PrefixRepetition,
+    #[value(name = "random-rerank")]
+    RandomRerank,
+}
+
+/// Ramp-up strategy for request rate.
+#[derive(clap::ValueEnum, Debug, Clone, Copy, PartialEq, Eq)]
+pub enum RampUpStrategy {
+    #[value(name = "linear")]
+    Linear,
+    #[value(name = "exponential")]
+    Exponential,
+}
+
+/// Strategy for assigning LoRA modules to requests.
+#[derive(clap::ValueEnum, Debug, Clone, Copy, PartialEq, Eq)]
+pub enum LoraAssignment {
+    #[value(name = "random")]
+    Random,
+    #[value(name = "round-robin")]
+    RoundRobin,
+}
+
+/// SPEED-Bench dataset split/config.
+#[derive(clap::ValueEnum, Debug, Clone, Copy, PartialEq, Eq)]
+pub enum SpeedBenchConfig {
+    #[value(name = "qualitative")]
+    Qualitative,
+    #[value(name = "throughput_1k")]
+    Throughput1k,
+    #[value(name = "throughput_2k")]
+    Throughput2k,
+    #[value(name = "throughput_8k")]
+    Throughput8k,
+    #[value(name = "throughput_16k")]
+    Throughput16k,
+    #[value(name = "throughput_32k")]
+    Throughput32k,
+}
+
+impl SpeedBenchConfig {
+    pub fn as_str(self) -> &'static str {
+        match self {
+            Self::Qualitative => "qualitative",
+            Self::Throughput1k => "throughput_1k",
+            Self::Throughput2k => "throughput_2k",
+            Self::Throughput8k => "throughput_8k",
+            Self::Throughput16k => "throughput_16k",
+            Self::Throughput32k => "throughput_32k",
+        }
+    }
+}
+
+impl fmt::Display for SpeedBenchConfig {
+    fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
+        f.write_str(self.as_str())
+    }
+}
+
+/// High-performance benchmark client for vLLM serving endpoints.
+#[derive(clap::Args, Debug, Clone)]
+pub struct BenchServeArgs {
+    /// The type of backend or endpoint to use for the benchmark.
+    #[arg(long, default_value = "openai")]
+    pub backend: BackendKind,
+
+    /// Server or API base url if not using http host and port.
+    #[arg(long)]
+    pub base_url: Option<String>,
+
+    /// Server host.
+    #[arg(long, default_value = "127.0.0.1")]
+    pub host: String,
+
+    /// Server port.
+    #[arg(long, default_value_t = 8000)]
+    pub port: u16,
+
+    /// API endpoint. Auto-selected based on --backend if not specified.
+    #[arg(long)]
+    pub endpoint: Option<String>,
+
+    /// Name of the model. If not specified, will fetch from server.
+    #[arg(long)]
+    pub model: Option<String>,
+
+    /// The model name used in the API (for --served-model-name).
+    #[arg(long)]
+    pub served_model_name: Option<String>,
+
+    /// Name or path of the tokenizer.
+    #[arg(long)]
+    pub tokenizer: Option<String>,
+
+    /// Tokenizer mode (auto, hf, slow, mistral).
+    #[arg(long, default_value = "auto")]
+    pub tokenizer_mode: String,
+
+    /// Skip initialization of tokenizer.
+    #[arg(long, default_value_t = false)]
+    pub skip_tokenizer_init: bool,
+
+    /// Trust remote code for tokenizer.
+    #[arg(long, default_value_t = false)]
+    pub trust_remote_code: bool,
+
+    /// Dataset name.
+    #[arg(long, default_value = "random")]
+    pub dataset_name: DatasetName,
+
+    /// General input length for datasets.
+    #[arg(long)]
+    pub input_len: Option<usize>,
+
+    /// General output length for datasets.
+    #[arg(long)]
+    pub output_len: Option<usize>,
+
+    /// Maximum model context length. Requests with prompt_len + output_len above this are filtered
+    /// out.
+    #[arg(long)]
+    pub max_model_len: Option<usize>,
+
+    /// Random dataset input length.
+    #[arg(long, default_value_t = 1024)]
+    pub random_input_len: usize,
+
+    /// Random dataset output length.
+    #[arg(long, default_value_t = 128)]
+    pub random_output_len: usize,
+
+    /// Random dataset prefix length.
+    #[arg(long, default_value_t = 0)]
+    pub random_prefix_len: usize,
+
+    /// Per-turn input length for turns 1+ in multi-turn mode.
+    /// 0 = fallback to --random-input-len for all turns.
+    /// Mirrors sglang bench_multiturn.py --sub-question-input-length.
+    #[arg(long, default_value_t = 0)]
+    pub per_turn_input_len: usize,
+
+    /// Range ratio for sampling input/output lengths, matching Python
+    /// `vllm bench serve`: lengths are drawn uniformly from
+    /// [len*(1-r), len*(1+r)]. 0.0 (the default) = exact target lengths.
+    /// Accepts a single float in [0, 1) or a JSON object
+    /// '{"input": r1, "output": r2}' for independent control.
+    /// NOTE: semantics changed — the old Rust-only form sampled [len*r, len]
+    /// with default 1.0; old values like 1.0 are now rejected.
+    #[arg(long, default_value = "0.0")]
+    pub random_range_ratio: String,
+
+    /// Batch multiple generated inputs into one request (embeddings/pooling
+    /// backends only). E.g. 8 sends "input": [t1..t8] per request. Mirrors
+    /// Python --random-batch-size. Default 1 = no batching.
+    #[arg(long, default_value_t = 1)]
+    pub random_batch_size: usize,
+
+    /// random-rerank: the served model is NOT a reranker (embedding-based
+    /// scoring). Changes query/document length accounting to mirror Python
+    /// --no-reranker.
+    #[arg(long, default_value_t = false)]
+    pub no_reranker: bool,
+
+    /// Bimodal prefix-cache (random dataset): fraction of prompts that are "warm"
+    /// and reuse a shared cached prefix. 0.0 = off (default). E.g. 0.8 = 80% warm
+    /// (prefix-cache hit), 20% cold (full prefill). Requires --random-cache-ratio > 0
+    /// and --prompt-token-ids. In this mode --random-input-len is the TOTAL length.
+    #[arg(long, default_value_t = 0.0)]
+    pub random_cache_hit_fraction: f64,
+
+    /// Bimodal prefix-cache (random dataset): fraction of each WARM prompt's length
+    /// that is the shared cached prefix. 0.0 = off (default). E.g. 0.95 = 95% cached,
+    /// 5% unique suffix. Used with --random-cache-hit-fraction.
+    #[arg(long, default_value_t = 0.0)]
+    pub random_cache_ratio: f64,
+
+    /// Send prompt as token ID arrays instead of text strings.
+    /// By default, prompts are decoded to text for maximum
+    /// compatibility. Enable this for pure vLLM deployments to skip server-side
+    /// tokenization (faster, exact token counts).
+    #[arg(long, default_value_t = false)]
+    pub prompt_token_ids: bool,
+
+    // --- Random multimodal dataset ---
+    /// Base number of multimodal items (images/videos) per request.
+    #[arg(long, default_value_t = 1)]
+    pub random_mm_base_items_per_request: usize,
+
+    /// Range ratio for varying the number of multimodal items per request.
+    /// Items sampled from [floor(n*(1-r)), ceil(n*(1+r))].
+    #[arg(long, default_value_t = 0.0)]
+    pub random_mm_num_mm_items_range_ratio: f64,
+
+    /// Per-modality hard caps as JSON, e.g. '{"image": 3, "video": 0}'.
+    #[arg(long, default_value = "{\"image\": 255, \"video\": 1}")]
+    pub random_mm_limit_mm_per_prompt: String,
+
+    /// Bucket config mapping (height,width,num_frames) to probability.
+    /// Uses Python-style syntax: '{(256,256,1): 0.5, (720,1280,1): 0.5}'.
+    /// num_frames=1 means image, num_frames>1 means video.
+    #[arg(long, default_value = "{(256,256,1): 0.5, (720,1280,1): 0.5}")]
+    pub random_mm_bucket_config: String,
+
+    /// Enable multimodal chat transformation for datasets that support it.
+    /// The dataset pre-builds the OpenAI chat `messages` array (text part +
+    /// multimodal items) at generation time, and the request sends it verbatim.
+    /// Mirrors Python's --enable-multimodal-chat. Currently applies to random-mm.
+    #[arg(long, default_value_t = false)]
+    pub enable_multimodal_chat: bool,
+
+    // --- Custom dataset (JSONL) ---
+    /// Output tokens per request for the custom dataset. Set to -1 to use the
+    /// per-line "output_tokens" field from the JSONL file instead.
+    #[arg(long, default_value_t = 256, allow_negative_numbers = true)]
+    pub custom_output_len: i64,
+
+    /// Skip applying a chat template to custom dataset prompts.
+    /// NOTE: the Rust client never renders chat templates client-side, so this
+    /// is always effectively on; passing it silences the informational notice.
+    #[arg(long, default_value_t = false)]
+    pub skip_chat_template: bool,
+
+    // --- Prefix repetition dataset ---
+    /// Shared-prefix token length for the prefix_repetition dataset.
+    #[arg(long, default_value_t = 256)]
+    pub prefix_repetition_prefix_len: usize,
+
+    /// Per-request random suffix token length for the prefix_repetition dataset.
+    #[arg(long, default_value_t = 256)]
+    pub prefix_repetition_suffix_len: usize,
+
+    /// Number of distinct shared prefixes for the prefix_repetition dataset.
+    /// Requests are split evenly across prefixes (num-prompts / num-prefixes each).
+    #[arg(long, default_value_t = 10)]
+    pub prefix_repetition_num_prefixes: usize,
+
+    /// Output tokens per request for the prefix_repetition dataset.
+    #[arg(long, default_value_t = 128)]
+    pub prefix_repetition_output_len: usize,
+
+    /// Number of prompts to generate.
+    #[arg(long, default_value_t = 1000)]
+    pub num_prompts: usize,
+
+    /// Number of requests per second. Use "inf" for all at once.
+    #[arg(long, default_value_t = f64::INFINITY)]
+    pub request_rate: f64,
+
+    /// Burstiness factor of request generation.
+    #[arg(long, default_value_t = 1.0)]
+    pub burstiness: f64,
+
+    /// Maximum number of concurrent requests.
+    #[arg(long)]
+    pub max_concurrency: Option<usize>,
+
+    /// Fraction of --max-concurrency at which the steady-state window opens.
+    /// Range: (0.0, 1.0]. Used only when --max-concurrency is set and
+    /// --request-rate is inf.
+    #[arg(long, default_value_t = 0.95)]
+    pub steady_state_threshold: f64,
+
+    /// Minimum steady-state window duration in seconds. Below this, a warning
+    /// is attached. If unset, computed as max(10.0, 0.1 * run_duration).
+    #[arg(long)]
+    pub steady_state_min_window: Option<f64>,
+
+    /// Disable steady-state metrics computation entirely.
+    #[arg(long, default_value_t = false)]
+    pub no_steady_state: bool,
+
+    /// Disable tqdm progress bar.
+    #[arg(long, default_value_t = false)]
+    pub disable_tqdm: bool,
+
+    /// Number of warmup requests.
+    #[arg(long, default_value_t = 0)]
+    pub num_warmups: usize,
+
+    /// Use vLLM profiling. --profiler-config must be provided on the server.
+    #[arg(long, default_value_t = false)]
+    pub profile: bool,
+
+    /// Minimum server batch size (num_requests_running) before starting the
+    /// profiler. When set, profiling is deferred until the /metrics endpoint
+    /// reports at least this many running requests, then captures for
+    /// --profile-duration seconds. Requires --profile.
+    #[arg(long)]
+    pub profile_batch_threshold: Option<usize>,
+
+    /// How many seconds to capture once the batch threshold is reached.
+    /// Defaults to 5. Requires --profile and --profile-batch-threshold.
+    #[arg(long, default_value_t = 5.0)]
+    pub profile_duration: f64,
+
+    /// Save benchmark results to a JSON file.
+    #[arg(long, default_value_t = false)]
+    pub save_result: bool,
+
+    /// Save detailed per-request results.
+    #[arg(long, default_value_t = false)]
+    pub save_detailed: bool,
+
+    /// Directory to save benchmark JSON results.
+    #[arg(long)]
+    pub result_dir: Option<String>,
+
+    /// Filename to save benchmark JSON results.
+    #[arg(long)]
+    pub result_filename: Option<String>,
+
+    /// Random seed.
+    #[arg(long, default_value_t = 0)]
+    pub seed: u64,
+
+    /// Set ignore_eos flag when sending the benchmark request.
+    #[arg(long, default_value_t = false)]
+    pub ignore_eos: bool,
+
+    /// Comma-separated list of metrics to report percentiles for.
+    #[arg(long)]
+    pub percentile_metrics: Option<String>,
+
+    /// Comma-separated list of percentiles for selected metrics.
+    #[arg(long, default_value = "99")]
+    pub metric_percentiles: String,
+
+    /// Comma-separated list of extra percentiles to show in sweep summaries.
+    #[arg(long)]
+    pub sweep_summary_percentiles: Option<String>,
+
+    /// The label (prefix) of the benchmark results.
+    #[arg(long)]
+    pub label: Option<String>,
+
+    /// Number of logprobs-per-token to compute.
+    #[arg(long)]
+    pub logprobs: Option<usize>,
+
+    /// Prefix for request IDs.
+    #[arg(long)]
+    pub request_id_prefix: Option<String>,
+
+    /// Maximum time to wait for endpoint readiness in seconds.
+    #[arg(long, default_value_t = 0)]
+    pub ready_check_timeout_sec: u64,
+
+    /// Key-value pairs for extra headers (KEY=VALUE).
+    #[arg(long = "header", num_args = 1..)]
+    pub headers: Option<Vec<String>>,
+
+    /// JSON string for extra body parameters.
+    #[arg(long)]
+    pub extra_body: Option<String>,
+
+    /// Key-value pairs for metadata (KEY=VALUE).
+    #[arg(long = "metadata", num_args = 1..)]
+    pub metadata: Option<Vec<String>>,
+
+    /// Dry run: only generate dataset and print stats, don't benchmark.
+    #[arg(long, default_value_t = false)]
+    pub dry_run: bool,
+
+    // --- Sampling parameters ---
+    /// Top-p sampling parameter. Only affects openai-compatible backends.
+    #[arg(long)]
+    pub top_p: Option<f64>,
+
+    /// Top-k sampling parameter. Only affects openai-compatible backends.
+    #[arg(long)]
+    pub top_k: Option<i64>,
+
+    /// Min-p sampling parameter. Only affects openai-compatible backends.
+    #[arg(long)]
+    pub min_p: Option<f64>,
+
+    /// Temperature sampling parameter. Only affects openai-compatible backends.
+    #[arg(long)]
+    pub temperature: Option<f64>,
+
+    /// Frequency penalty sampling parameter. Only affects openai-compatible backends.
+    #[arg(long)]
+    pub frequency_penalty: Option<f64>,
+
+    /// Presence penalty sampling parameter. Only affects openai-compatible backends.
+    #[arg(long)]
+    pub presence_penalty: Option<f64>,
+
+    /// Repetition penalty sampling parameter. Only affects openai-compatible backends.
+    #[arg(long)]
+    pub repetition_penalty: Option<f64>,
+
+    // --- SSL ---
+    /// Disable SSL certificate verification.
+    #[arg(long, default_value_t = false)]
+    pub insecure: bool,
+
+    // --- Ramp-up ---
+    /// Ramp-up strategy for request rate (linear or exponential).
+    #[arg(long)]
+    pub ramp_up_strategy: Option<RampUpStrategy>,
+
+    /// Starting request rate for ramp-up (RPS).
+    #[arg(long)]
+    pub ramp_up_start_rps: Option<f64>,
+
+    /// Ending request rate for ramp-up (RPS).
+    #[arg(long)]
+    pub ramp_up_end_rps: Option<f64>,
+
+    // --- Goodput ---
+    /// Service level objectives for goodput as "KEY:VALUE" pairs (e.g. ttft:100 tpot:50 e2el:500).
+    /// Values are in milliseconds.
+    #[arg(long = "goodput", num_args = 1..)]
+    pub goodput: Option<Vec<String>>,
+
+    // --- Result ---
+    /// Append the benchmark result to the existing JSON file.
+    #[arg(long, default_value_t = false)]
+    pub append_result: bool,
+
+    // --- ShareGPT dataset ---
+    /// Path to dataset file (required for sharegpt dataset).
+    #[arg(long)]
+    pub dataset_path: Option<String>,
+
+    /// Override output length for ShareGPT dataset.
+    #[arg(long)]
+    pub sharegpt_output_len: Option<usize>,
+
+    /// Do not oversample if dataset is smaller than num_prompts.
+    #[arg(long, default_value_t = false)]
+    pub no_oversample: bool,
+
+    /// Do not shuffle the dataset.
+    #[arg(long, default_value_t = false)]
+    pub disable_shuffle: bool,
+
+    // --- Sonnet dataset ---
+    /// Number of input tokens per request (sonnet dataset).
+    #[arg(long, default_value_t = crate::datasets::sonnet::DEFAULT_INPUT_LEN)]
+    pub sonnet_input_len: usize,
+
+    /// Number of output tokens per request (sonnet dataset).
+    #[arg(long, default_value_t = crate::datasets::sonnet::DEFAULT_OUTPUT_LEN)]
+    pub sonnet_output_len: usize,
+
+    /// Number of prefix tokens shared across requests (sonnet dataset).
+    #[arg(long, default_value_t = crate::datasets::sonnet::DEFAULT_PREFIX_LEN)]
+    pub sonnet_prefix_len: usize,
+
+    /// SPEED-Bench config/split (qualitative, throughput_1k, throughput_2k, throughput_8k,
+    /// throughput_16k, throughput_32k).
+    #[arg(long, default_value = "qualitative")]
+    pub speed_bench_config: SpeedBenchConfig,
+
+    /// Filter SPEED-Bench by category (e.g. low_entropy, high_entropy, coding, math).
+    #[arg(long)]
+    pub speed_bench_category: Option<String>,
+
+    /// Truncate SPEED-Bench prompts to at most this many tokens.
+    /// Useful for creating custom input lengths from larger splits (e.g. --speed-bench-config
+    /// throughput_16k --speed-bench-max-input-len 10240).
+    #[arg(long)]
+    pub speed_bench_max_input_len: Option<usize>,
+
+    // --- HuggingFace dataset ---
+    /// HuggingFace dataset split (e.g. train, test, validation).
+    #[arg(long)]
+    pub hf_split: Option<String>,
+
+    /// HuggingFace dataset subset/config name.
+    #[arg(long)]
+    pub hf_subset: Option<String>,
+
+    /// Fixed output length for HF dataset requests (overrides dataset-derived length).
+    #[arg(long)]
+    pub hf_output_len: Option<usize>,
+
+    /// Column name containing the prompt text. Auto-detected if not specified.
+    #[arg(long)]
+    pub hf_text_column: Option<String>,
+
+    // --- Compare mode ---
+    /// Compare two benchmark result JSON files (e.g. --compare a.json b.json).
+    /// Prints side-by-side metrics with delta and % change. Skips benchmarking.
+    #[arg(long = "compare", num_args = 2, value_names = ["FILE_A", "FILE_B"])]
+    pub compare: Option<Vec<String>>,
+
+    // --- Sweep mode ---
+    /// Sweep over max-concurrency values (comma-separated, e.g. --sweep-max-concurrency
+    /// 1,10,50,100,500).
+    #[arg(long)]
+    pub sweep_max_concurrency: Option<String>,
+
+    /// When sweeping concurrency, set num_prompts = concurrency * this factor for each sweep
+    /// point.
+    #[arg(long)]
+    pub sweep_num_prompts_factor: Option<usize>,
+
+    /// Sweep over request-rate values (comma-separated, supports "inf", e.g. --sweep-request-rate
+    /// 1,10,100,inf).
+    #[arg(long)]
+    pub sweep_request_rate: Option<String>,
+
+    /// Reset the server's prefix cache before each sweep iteration.
+    /// Requires VLLM_SERVER_DEV_MODE=1 on the vLLM server.
+    #[arg(long, default_value_t = false)]
+    pub reset_prefix_cache: bool,
+
+    // --- Multi-run ---
+    /// Number of benchmark runs for statistical aggregation.
+    #[arg(long, default_value_t = 1)]
+    pub num_runs: usize,
+
+    // --- Multi-turn conversation benchmark ---
+    /// Enable multi-turn conversation benchmark mode.
+    #[arg(long, default_value_t = false)]
+    pub multi_turn: bool,
+
+    /// Number of turns per conversation in synthetic multi-turn mode.
+    #[arg(long, default_value_t = 3)]
+    pub multi_turn_num_turns: usize,
+
+    /// Minimum turns per conversation. 0 = use --multi-turn-num-turns.
+    #[arg(long, default_value_t = 0)]
+    pub multi_turn_min_turns: usize,
+
+    /// Maximum turns per conversation.
+    /// For synthetic multi-turn, 0 = use --multi-turn-num-turns.
+    /// For ShareGPT multi-turn, 0 = uncapped.
+    #[arg(long, default_value_t = 0)]
+    pub multi_turn_max_turns: usize,
+
+    /// Number of concurrent conversations (defaults to max-concurrency or num-prompts).
+    #[arg(long)]
+    pub multi_turn_concurrency: Option<usize>,
+
+    /// Delay between turns in milliseconds (simulates user think time).
+    #[arg(long, default_value_t = 0)]
+    pub multi_turn_delay_ms: u64,
+
+    /// Fraction of per-turn input tokens shared across ALL conversations (0.0–1.0).
+    /// When > 0, enables prefix sharing mode: each turn sends a fixed-length message
+    /// (no history accumulation). Only works with --dataset-name random.
+    #[arg(long, default_value_t = 0.0)]
+    pub multi_turn_prefix_global_ratio: f64,
+
+    /// Fraction of per-turn input tokens shared within each conversation (0.0–1.0).
+    /// When > 0, enables prefix sharing mode: each turn sends a fixed-length message
+    /// (no history accumulation). Only works with --dataset-name random.
+    #[arg(long, default_value_t = 0.0)]
+    pub multi_turn_prefix_conversation_ratio: f64,
+
+    // --- LoRA ---
+    /// LoRA adapter names registered on the server (server-side
+    /// `--lora-modules name=path`). Each request's `model` field is rewritten
+    /// to one of these names; tokenizer and other endpoints keep using --model.
+    /// In multi-turn mode, one adapter is assigned per conversation (sticky
+    /// across turns).
+    #[arg(long = "lora-modules", num_args = 1..)]
+    pub lora_modules: Option<Vec<String>>,
+
+    /// Strategy for assigning LoRA adapters to requests.
+    /// 'random' (default) picks uniformly at random; 'round-robin' cycles
+    /// through `--lora-modules` deterministically (i % N).
+    #[arg(long = "lora-assignment", default_value = "random")]
+    pub lora_assignment: LoraAssignment,
+}
+
+impl BenchServeArgs {
+    /// Resolve the base URL from explicit --base-url or from --host/--port.
+    pub fn resolve_base_url(&self) -> String {
+        if let Some(ref base) = self.base_url {
+            base.clone()
+        } else {
+            format!("http://{}:{}", self.host, self.port)
+        }
+    }
+
+    /// Resolve the API endpoint, auto-selecting based on backend if not explicit.
+    pub fn resolve_endpoint(&self) -> String {
+        if let Some(ref ep) = self.endpoint {
+            return ep.clone();
+        }
+        match self.backend {
+            BackendKind::OpenaiChat => "/v1/chat/completions".to_string(),
+            BackendKind::Vllm | BackendKind::Openai => "/v1/completions".to_string(),
+            BackendKind::OpenaiEmbeddings | BackendKind::OpenaiEmbeddingsChat => {
+                "/v1/embeddings".to_string()
+            }
+            BackendKind::VllmPooling => "/v1/pooling".to_string(),
+            BackendKind::VllmRerank => "/v1/rerank".to_string(),
+        }
+    }
+
+    /// Resolve the full API URL.
+    pub fn resolve_api_url(&self) -> String {
+        format!("{}{}", self.resolve_base_url(), self.resolve_endpoint())
+    }
+
+    /// Parse extra headers from KEY=VALUE pairs.
+    pub fn parse_headers(
+        &self,
+    ) -> crate::error::Result<Option<std::collections::HashMap<String, String>>> {
+        match &self.headers {
+            None => Ok(None),
+            Some(items) => {
+                let mut map = std::collections::HashMap::new();
+                for item in items {
+                    let (k, v) = item.split_once('=').ok_or_else(|| {
+                        crate::error::BenchError::Config(
+                            "Invalid header format. Use KEY=VALUE".into(),
+                        )
+                    })?;
+                    map.insert(k.trim().to_string(), v.trim().to_string());
+                }
+                Ok(Some(map))
+            }
+        }
+    }
+
+    /// Parse extra body JSON.
+    pub fn parse_extra_body(&self) -> crate::error::Result<Option<serde_json::Value>> {
+        match &self.extra_body {
+            None => Ok(None),
+            Some(s) => {
+                let v: serde_json::Value = serde_json::from_str(s).map_err(|e| {
+                    crate::error::BenchError::Config(format!("Invalid --extra-body JSON: {e}"))
+                })?;
+                Ok(Some(v))
+            }
+        }
+    }
+
+    /// Generate the request ID prefix (auto-generate if not provided).
+    pub fn get_request_id_prefix(&self) -> String {
+        self.request_id_prefix
+            .clone()
+            .unwrap_or_else(|| format!("bench-{}-", &uuid::Uuid::new_v4().to_string()[..8]))
+    }
+
+    /// Resolve input/output lengths, applying --input-len/--output-len overrides.
+    pub fn resolved_random_input_len(&self) -> usize {
+        self.input_len.unwrap_or(self.random_input_len)
+    }
+
+    pub fn resolved_random_output_len(&self) -> usize {
+        self.output_len.unwrap_or(self.random_output_len)
+    }
+
+    pub fn resolved_per_turn_input_len(&self) -> usize {
+        if self.per_turn_input_len > 0 {
+            self.per_turn_input_len
+        } else {
+            self.resolved_random_input_len()
+        }
+    }
+}
@@ -0,0 +1,302 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+use crate::error::{BenchError, Result};
+
+/// Metric definition for comparison: name, JSON key, and whether lower is better.
+struct MetricDef {
+    label: &'static str,
+    key: &'static str,
+    lower_is_better: bool,
+}
+
+const METRICS: &[MetricDef] = &[
+    MetricDef {
+        label: "Request throughput (req/s)",
+        key: "request_throughput",
+        lower_is_better: false,
+    },
+    MetricDef {
+        label: "Output throughput (tok/s)",
+        key: "output_throughput",
+        lower_is_better: false,
+    },
+    MetricDef {
+        label: "Total token throughput (tok/s)",
+        key: "total_token_throughput",
+        lower_is_better: false,
+    },
+    MetricDef {
+        label: "Peak output tokens/s",
+        key: "max_output_tokens_per_s",
+        lower_is_better: false,
+    },
+    MetricDef {
+        label: "Peak concurrent requests",
+        key: "max_concurrent_requests",
+        lower_is_better: false,
+    },
+    MetricDef {
+        label: "Mean TTFT (ms)",
+        key: "mean_ttft_ms",
+        lower_is_better: true,
+    },
+    MetricDef {
+        label: "Median TTFT (ms)",
+        key: "median_ttft_ms",
+        lower_is_better: true,
+    },
+    MetricDef {
+        label: "P99 TTFT (ms)",
+        key: "p99_ttft_ms",
+        lower_is_better: true,
+    },
+    MetricDef {
+        label: "Mean TPOT (ms)",
+        key: "mean_tpot_ms",
+        lower_is_better: true,
+    },
+    MetricDef {
+        label: "Median TPOT (ms)",
+        key: "median_tpot_ms",
+        lower_is_better: true,
+    },
+    MetricDef {
+        label: "P99 TPOT (ms)",
+        key: "p99_tpot_ms",
+        lower_is_better: true,
+    },
+    MetricDef {
+        label: "Mean ITL (ms)",
+        key: "mean_itl_ms",
+        lower_is_better: true,
+    },
+    MetricDef {
+        label: "Median ITL (ms)",
+        key: "median_itl_ms",
+        lower_is_better: true,
+    },
+    MetricDef {
+        label: "P99 ITL (ms)",
+        key: "p99_itl_ms",
+        lower_is_better: true,
+    },
+    MetricDef {
+        label: "Mean E2EL (ms)",
+        key: "mean_e2el_ms",
+        lower_is_better: true,
+    },
+    MetricDef {
+        label: "Median E2EL (ms)",
+        key: "median_e2el_ms",
+        lower_is_better: true,
+    },
+    MetricDef {
+        label: "P99 E2EL (ms)",
+        key: "p99_e2el_ms",
+        lower_is_better: true,
+    },
+    MetricDef {
+        label: "Completed requests",
+        key: "completed",
+        lower_is_better: false,
+    },
+    MetricDef {
+        label: "Failed requests",
+        key: "failed",
+        lower_is_better: true,
+    },
+    MetricDef {
+        label: "Duration (s)",
+        key: "duration",
+        lower_is_better: true,
+    },
+];
+
+const STEADY_STATE_METRICS: &[MetricDef] = &[
+    MetricDef {
+        label: "SS Request throughput (req/s)",
+        key: "request_throughput",
+        lower_is_better: false,
+    },
+    MetricDef {
+        label: "SS Output throughput (tok/s)",
+        key: "output_throughput",
+        lower_is_better: false,
+    },
+    MetricDef {
+        label: "SS Input throughput (tok/s)",
+        key: "input_throughput",
+        lower_is_better: false,
+    },
+    MetricDef {
+        label: "SS Total token throughput (tok/s)",
+        key: "total_token_throughput",
+        lower_is_better: false,
+    },
+    MetricDef {
+        label: "SS Mean TTFT (ms)",
+        key: "mean_ttft_ms",
+        lower_is_better: true,
+    },
+    MetricDef {
+        label: "SS Median TTFT (ms)",
+        key: "median_ttft_ms",
+        lower_is_better: true,
+    },
+    MetricDef {
+        label: "SS Mean TPOT (ms)",
+        key: "mean_tpot_ms",
+        lower_is_better: true,
+    },
+    MetricDef {
+        label: "SS Median TPOT (ms)",
+        key: "median_tpot_ms",
+        lower_is_better: true,
+    },
+    MetricDef {
+        label: "SS P90 TPOT (ms)",
+        key: "p90_tpot_ms",
+        lower_is_better: true,
+    },
+    MetricDef {
+        label: "SS P99 TPOT (ms)",
+        key: "p99_tpot_ms",
+        lower_is_better: true,
+    },
+];
+
+/// Compare two benchmark result JSON files and print a side-by-side table.
+pub fn compare_results(file_a: &str, file_b: &str) -> Result<()> {
+    let json_a = load_result_json(file_a)?;
+    let json_b = load_result_json(file_b)?;
+
+    // Print header with file context
+    let model_a = json_a.get("model_id").and_then(|v| v.as_str()).unwrap_or("?");
+    let model_b = json_b.get("model_id").and_then(|v| v.as_str()).unwrap_or("?");
+    let date_a = json_a.get("date").and_then(|v| v.as_str()).unwrap_or("?");
+    let date_b = json_b.get("date").and_then(|v| v.as_str()).unwrap_or("?");
+
+    println!("{:=^90}", " Benchmark Comparison ");
+    println!("  A: {} (model: {}, date: {})", file_a, model_a, date_a);
+    println!("  B: {} (model: {}, date: {})", file_b, model_b, date_b);
+    println!();
+
+    // Print comparison table
+    println!(
+        "{:<35} {:>12} {:>12} {:>10} {:>8}",
+        "Metric", "A", "B", "Delta", "Change"
+    );
+    println!("{:-<35} {:->12} {:->12} {:->10} {:->8}", "", "", "", "", "");
+
+    for metric in METRICS {
+        let val_a = get_f64(&json_a, metric.key);
+        let val_b = get_f64(&json_b, metric.key);
+
+        match (val_a, val_b) {
+            (Some(a), Some(b)) => print_diff_row(metric, a, b),
+            _ => {
+                // One or both values missing — skip
+            }
+        }
+    }
+
+    // Steady-state section — both sides must have the block; otherwise render N/A.
+    let ss_a = json_a.get("steady_state");
+    let ss_b = json_b.get("steady_state");
+    let both_present =
+        matches!(ss_a, Some(v) if !v.is_null()) && matches!(ss_b, Some(v) if !v.is_null());
+
+    println!();
+    println!("{:=^70}", " Steady-State Comparison ");
+    if !both_present {
+        println!("N/A — one or both runs have no steady-state window");
+    } else {
+        let ss_a = ss_a.unwrap();
+        let ss_b = ss_b.unwrap();
+        for m in STEADY_STATE_METRICS {
+            let a = ss_a.get(m.key).and_then(|v| v.as_f64());
+            let b = ss_b.get(m.key).and_then(|v| v.as_f64());
+            match (a, b) {
+                (Some(a), Some(b)) => print_diff_row(m, a, b),
+                _ => println!("{:<35} N/A", m.label),
+            }
+        }
+    }
+
+    println!("{:=<90}", "");
+    println!();
+    println!("Legend: + = improvement, - = regression (relative to A → B)");
+
+    Ok(())
+}
+
+fn print_diff_row(metric: &MetricDef, a: f64, b: f64) {
+    let delta = b - a;
+    let pct = if a.abs() > 1e-10 {
+        (delta / a) * 100.0
+    } else if b.abs() > 1e-10 {
+        f64::INFINITY
+    } else {
+        0.0
+    };
+
+    // Determine if change is good/bad/neutral
+    let marker = if delta.abs() < 1e-10 {
+        " "
+    } else if metric.lower_is_better {
+        if delta < 0.0 { "+" } else { "-" }
+    } else if delta > 0.0 {
+        "+"
+    } else {
+        "-"
+    };
+
+    let delta_str = format_delta(delta);
+    let pct_str = if pct.is_infinite() {
+        "inf%".to_string()
+    } else {
+        format!("{:+.1}%", pct)
+    };
+
+    println!(
+        "{:<35} {:>12} {:>12} {:>10} {:>7}{}",
+        metric.label,
+        format_value(a),
+        format_value(b),
+        delta_str,
+        pct_str,
+        marker,
+    );
+}
+
+fn load_result_json(path: &str) -> Result<serde_json::Value> {
+    let content = std::fs::read_to_string(path)
+        .map_err(|e| BenchError::Config(format!("Cannot read result file '{path}': {e}")))?;
+
+    // Support JSONL: take the last line (most recent run)
+    let json_str = content.lines().rfind(|l| !l.trim().is_empty()).unwrap_or(&content);
+
+    serde_json::from_str(json_str)
+        .map_err(|e| BenchError::Config(format!("Cannot parse JSON from '{path}': {e}")))
+}
+
+fn get_f64(json: &serde_json::Value, key: &str) -> Option<f64> {
+    json.get(key).and_then(|v| v.as_f64())
+}
+
+fn format_value(v: f64) -> String {
+    if v == v.floor() && v.abs() < 1e12 {
+        format!("{}", v as i64)
+    } else {
+        format!("{:.2}", v)
+    }
+}
+
+fn format_delta(d: f64) -> String {
+    if d == d.floor() && d.abs() < 1e12 {
+        format!("{:+}", d as i64)
+    } else {
+        format!("{:+.2}", d)
+    }
+}
@@ -0,0 +1,181 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+//! Custom dataset: JSONL file with one request per line.
+//!
+//! ```jsonl
+//! {"prompt": "What is the capital of India?", "output_tokens": 10}
+//! {"prompt": "What is the capital of Iran?", "output_tokens": 1520}
+//! ```
+//!
+//! Mirrors Python's `CustomDataset`. `output_tokens` is optional unless
+//! `--custom-output-len -1` is passed. Unlike Python, prompts are always sent
+//! raw (no client-side chat template; see `--skip-chat-template`).
+
+use std::sync::Arc;
+
+use rand::SeedableRng;
+use rand::rngs::StdRng;
+use rand::seq::SliceRandom;
+use serde::Deserialize;
+
+use super::SampleRequest;
+use crate::error::{BenchError, Result};
+use crate::tokenizer::TokenizerKind;
+
+#[derive(Deserialize)]
+struct CustomLine {
+    prompt: String,
+    output_tokens: Option<serde_json::Value>,
+}
+
+/// Load the custom JSONL dataset.
+///
+/// `output_len < 0` means "use the per-line output_tokens field" (Python's
+/// `--custom-output-len -1`); otherwise `output_len` applies to every request.
+pub fn load_custom_dataset(
+    tokenizer: &TokenizerKind,
+    path: &str,
+    num_requests: usize,
+    output_len: i64,
+    seed: u64,
+    request_id_prefix: &str,
+    no_oversample: bool,
+    disable_shuffle: bool,
+) -> Result<Vec<SampleRequest>> {
+    let content = std::fs::read_to_string(path)
+        .map_err(|e| BenchError::Config(format!("Failed to read custom dataset '{path}': {e}")))?;
+
+    let mut lines: Vec<CustomLine> = Vec::new();
+    for (lineno, line) in content.lines().enumerate() {
+        let line = line.trim();
+        if line.is_empty() {
+            continue;
+        }
+        let parsed: CustomLine = serde_json::from_str(line).map_err(|e| {
+            BenchError::Config(format!(
+                "Invalid JSONL at {path}:{}: {e} (each line must be an object \
+                 with a 'prompt' field)",
+                lineno + 1
+            ))
+        })?;
+        lines.push(parsed);
+    }
+    if lines.is_empty() {
+        return Err(BenchError::Config(format!(
+            "Custom dataset '{path}' contains no entries"
+        )));
+    }
+
+    // Python shuffles the loaded data (seeded) before taking num_requests.
+    if !disable_shuffle {
+        let mut rng = StdRng::seed_from_u64(seed);
+        lines.shuffle(&mut rng);
+    }
+
+    let mut requests: Vec<SampleRequest> = Vec::with_capacity(num_requests.min(lines.len()));
+    for (i, item) in lines.iter().enumerate() {
+        if requests.len() >= num_requests {
+            break;
+        }
+
+        let expected_output_len = if output_len < 0 {
+            let raw = item.output_tokens.as_ref().ok_or_else(|| {
+                BenchError::Config(
+                    "custom dataset: --custom-output-len -1 requires an \
+                     'output_tokens' field on every line"
+                        .into(),
+                )
+            })?;
+            raw.as_i64().filter(|v| *v > 0).ok_or_else(|| {
+                BenchError::Config(format!(
+                    "custom dataset: invalid 'output_tokens' value {raw}: \
+                         must be a positive integer"
+                ))
+            })? as usize
+        } else {
+            output_len as usize
+        };
+
+        let prompt_len = tokenizer.encode(&item.prompt, true)?.len();
+        requests.push(SampleRequest {
+            prompt: Arc::from(item.prompt.as_str()),
+            prompt_len,
+            expected_output_len,
+            request_id: Some(format!("{request_id_prefix}{i}")),
+            ..Default::default()
+        });
+    }
+
+    super::oversample_requests(
+        &mut requests,
+        num_requests,
+        request_id_prefix,
+        no_oversample,
+    );
+    Ok(requests)
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+
+    fn write_temp_jsonl(name: &str, content: &str) -> String {
+        let path = std::env::temp_dir().join(format!("vllm-bench-custom-{name}.jsonl"));
+        std::fs::write(&path, content).unwrap();
+        path.to_string_lossy().into_owned()
+    }
+
+    /// gpt2 via built-in tiktoken encoding — loads without network access.
+    fn test_tokenizer() -> TokenizerKind {
+        crate::tokenizer::load_tokenizer("gpt2", false, None)
+            .expect("gpt2 built-in tiktoken should always load without network")
+    }
+
+    #[test]
+    fn test_load_custom_dataset_basic() {
+        let path = write_temp_jsonl(
+            "basic",
+            r#"{"prompt": "hello world", "output_tokens": 10}
+{"prompt": "foo bar baz", "output_tokens": 20}
+"#,
+        );
+        let reqs = load_custom_dataset(&test_tokenizer(), &path, 2, 256, 0, "t-", true, true)
+            .expect("load should succeed");
+        assert_eq!(reqs.len(), 2);
+        // Fixed output_len (256) wins over per-line output_tokens by default
+        assert!(reqs.iter().all(|r| r.expected_output_len == 256));
+        assert_eq!(&*reqs[0].prompt, "hello world");
+        assert!(reqs[0].prompt_len > 0);
+    }
+
+    #[test]
+    fn test_load_custom_dataset_per_line_output_tokens() {
+        let path = write_temp_jsonl(
+            "perline",
+            r#"{"prompt": "hello", "output_tokens": 10}
+{"prompt": "world", "output_tokens": 20}
+"#,
+        );
+        let reqs = load_custom_dataset(&test_tokenizer(), &path, 2, -1, 0, "t-", true, true)
+            .expect("load should succeed");
+        assert_eq!(reqs[0].expected_output_len, 10);
+        assert_eq!(reqs[1].expected_output_len, 20);
+    }
+
+    #[test]
+    fn test_load_custom_dataset_missing_output_tokens_errors() {
+        let path = write_temp_jsonl("missing", r#"{"prompt": "hello"}"#);
+        let err = load_custom_dataset(&test_tokenizer(), &path, 1, -1, 0, "t-", true, true)
+            .expect_err("should fail without output_tokens");
+        assert!(err.to_string().contains("output_tokens"));
+    }
+
+    #[test]
+    fn test_load_custom_dataset_missing_prompt_errors() {
+        let path = write_temp_jsonl("noprompt", r#"{"text": "hello"}"#);
+        assert!(
+            load_custom_dataset(&test_tokenizer(), &path, 1, 256, 0, "t-", true, true).is_err()
+        );
+    }
+}
@@ -0,0 +1,207 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+pub mod custom;
+pub mod hf_dataset;
+pub mod multi_turn;
+pub mod prefix_repetition;
+pub mod random;
+pub mod random_mm;
+pub mod random_rerank;
+pub mod sharegpt;
+pub mod sonnet;
+pub mod speed_bench;
+
+use std::sync::Arc;
+
+/// Represents a single inference request for benchmarking.
+/// Matches Python's SampleRequest dataclass from datasets.py:71-82.
+///
+/// `prompt` uses `Arc<str>` to avoid expensive String clones when distributing
+/// requests across tokio tasks. At 100k prompts with 8k tokens each, this saves
+/// ~3GB of peak memory vs cloning String per task.
+#[derive(Debug, Clone)]
+pub struct SampleRequest {
+    pub prompt: Arc<str>,
+    pub prompt_len: usize,
+    pub expected_output_len: usize,
+    pub request_id: Option<String>,
+    /// Pre-computed token IDs for this prompt.
+    /// When set, the completions backend sends these directly via `prompt_token_ids`
+    /// instead of the text `prompt`, avoiding server-side re-tokenization.
+    pub prompt_token_ids: Option<Arc<[u32]>>,
+    /// Multimodal content items as pre-serialized JSON fragments.
+    /// Each `Arc<str>` is a complete JSON object string, e.g.
+    /// `{"type":"image_url","image_url":{"url":"data:image/jpeg;base64,..."}}`
+    ///
+    /// Pre-serialized to avoid:
+    /// 1. `serde_json::Value` tree overhead (3 Maps + keys per image)
+    /// 2. Deep-cloning ~200KB+ base64 data when building request payloads
+    ///
+    /// Double-`Arc` for zero-cost sharing: outer Arc for the slice, inner Arc for each fragment.
+    pub multi_modal_content: Option<Arc<[Arc<str>]>>,
+    /// Pre-serialized OpenAI chat `messages` array as a complete JSON string,
+    /// e.g. `[{"role":"user","content":[{"type":"text","text":"..."},{"type":"image_url",...}]}]`.
+    ///
+    /// Set by datasets when `--enable-multimodal-chat` is on (mirrors Python's
+    /// `apply_multimodal_chat_transformation`: the dataset builds the chat messages
+    /// and the backend sends them verbatim). When set, `multi_modal_content` is None
+    /// and the mm items are embedded here instead. `prompt` still holds the text part
+    /// for token accounting and /tokenize verification.
+    pub chat_messages_json: Option<Arc<str>>,
+    /// Multiple text inputs for one request (pooling backends only).
+    /// Embeddings send it as `"input": [t1, t2, ...]` (--random-batch-size);
+    /// rerank sends `[0]` as the query and `[1..]` as documents (random-rerank).
+    /// Mirrors Python's list-valued `SampleRequest.prompt`.
+    pub prompt_list: Option<Arc<[Arc<str>]>>,
+}
+
+impl Default for SampleRequest {
+    /// Empty request; struct-update base so dataset builders only spell out the
+    /// fields they set (new optional fields then don't touch every call site).
+    fn default() -> Self {
+        Self {
+            prompt: Arc::from(""),
+            prompt_len: 0,
+            expected_output_len: 0,
+            request_id: None,
+            prompt_token_ids: None,
+            multi_modal_content: None,
+            chat_messages_json: None,
+            prompt_list: None,
+        }
+    }
+}
+
+/// Oversample `requests` up to `num_requests` by cloning random entries
+/// (seeded by list length for determinism), renumbering their request ids.
+/// No-op when enough samples exist, `no_oversample` is set, or the list is empty.
+/// Mirrors Python `BenchmarkDataset.maybe_oversample_requests`.
+pub fn oversample_requests(
+    requests: &mut Vec<SampleRequest>,
+    num_requests: usize,
+    request_id_prefix: &str,
+    no_oversample: bool,
+) {
+    use rand::rngs::StdRng;
+    use rand::{Rng, SeedableRng};
+
+    if requests.len() >= num_requests || requests.is_empty() {
+        return;
+    }
+    if no_oversample {
+        println!(
+            "Skipping oversampling. Total samples: {} (requested: {num_requests})",
+            requests.len()
+        );
+        return;
+    }
+    let original_len = requests.len();
+    let mut rng = StdRng::seed_from_u64(original_len as u64);
+    for i in 0..(num_requests - original_len) {
+        let mut req = requests[rng.random_range(0..original_len)].clone();
+        req.request_id = Some(format!("{request_id_prefix}{}", original_len + i));
+        requests.push(req);
+    }
+    println!(
+        "Oversampled requests from {original_len} to {} total samples.",
+        requests.len()
+    );
+}
+
+/// Group already-generated single-input requests into batched requests of
+/// `batch_size` inputs each (embeddings/pooling only). Mirrors Python
+/// `RandomDataset.sample` batching: prompt becomes a list, prompt_len is the
+/// sum over the batch, request ids are renumbered per batch.
+/// `batch_size <= 1` returns the input unchanged.
+pub fn batch_requests(
+    requests: Vec<SampleRequest>,
+    batch_size: usize,
+    request_id_prefix: &str,
+) -> Vec<SampleRequest> {
+    if batch_size <= 1 {
+        return requests;
+    }
+    requests
+        .chunks(batch_size)
+        .enumerate()
+        .map(|(batch_idx, batch)| SampleRequest {
+            prompt_list: Some(batch.iter().map(|r| r.prompt.clone()).collect()),
+            prompt_len: batch.iter().map(|r| r.prompt_len).sum(),
+            expected_output_len: 0,
+            request_id: Some(format!("{request_id_prefix}{batch_idx}")),
+            ..Default::default()
+        })
+        .collect()
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+
+    fn req(prompt: &str, len: usize) -> SampleRequest {
+        SampleRequest {
+            prompt: Arc::from(prompt),
+            prompt_len: len,
+            expected_output_len: 128,
+            ..Default::default()
+        }
+    }
+
+    #[test]
+    fn test_batch_requests_groups_and_sums() {
+        let reqs = vec![
+            req("a", 10),
+            req("b", 20),
+            req("c", 30),
+            req("d", 40),
+            req("e", 50),
+        ];
+        let batched = batch_requests(reqs, 2, "t-");
+        assert_eq!(batched.len(), 3); // 2 + 2 + 1
+        let first = batched[0].prompt_list.as_ref().unwrap();
+        assert_eq!(first.len(), 2);
+        assert_eq!(&*first[0], "a");
+        assert_eq!(batched[0].prompt_len, 30);
+        assert_eq!(batched[0].expected_output_len, 0);
+        assert_eq!(batched[0].request_id.as_deref(), Some("t-0"));
+        assert_eq!(batched[2].prompt_list.as_ref().unwrap().len(), 1);
+        assert_eq!(batched[2].prompt_len, 50);
+    }
+
+    #[test]
+    fn test_batch_requests_size_one_is_identity() {
+        let reqs = vec![req("a", 10), req("b", 20)];
+        let out = batch_requests(reqs, 1, "t-");
+        assert_eq!(out.len(), 2);
+        assert!(out[0].prompt_list.is_none());
+        assert_eq!(&*out[0].prompt, "a");
+    }
+
+    #[test]
+    fn test_oversample_requests() {
+        let mut reqs = vec![req("a", 10), req("b", 20)];
+        oversample_requests(&mut reqs, 5, "t-", false);
+        assert_eq!(reqs.len(), 5);
+        assert_eq!(reqs[4].request_id.as_deref(), Some("t-4"));
+
+        let mut reqs = vec![req("a", 10)];
+        oversample_requests(&mut reqs, 5, "t-", true); // no_oversample
+        assert_eq!(reqs.len(), 1);
+    }
+}
+
+/// A single turn in a multi-turn conversation.
+#[derive(Debug, Clone)]
+pub struct ConversationTurn {
+    pub user_message: Arc<str>,
+    pub user_message_len: usize,
+    pub expected_output_len: usize,
+}
+
+/// A complete multi-turn conversation with all turns pre-generated.
+#[derive(Debug, Clone)]
+pub struct MultiTurnConversation {
+    pub conversation_id: String,
+    pub turns: Vec<ConversationTurn>,
+}
@@ -0,0 +1,777 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+use std::sync::Arc;
+
+use rand::rngs::StdRng;
+use rand::seq::SliceRandom;
+use rand::{Rng, SeedableRng};
+use rayon::prelude::*;
+
+use super::{ConversationTurn, MultiTurnConversation};
+use crate::error::{BenchError, Result};
+use crate::tokenizer::TokenizerKind;
+
+/// Configuration for generating random multi-turn conversations.
+#[derive(Debug, Clone)]
+pub struct MultiTurnRandomConfig {
+    pub num_conversations: usize,
+    pub min_turns: usize,
+    pub max_turns: usize,
+    /// Shared prefix length prepended to the conversation.
+    ///
+    /// In normal accumulated-history mode this is added to turn 0, so all
+    /// later turns inherit it through history. In no-history prefix-sharing
+    /// mode it is added to every independent turn.
+    pub prefix_len: usize,
+    /// Input length for turn 0.
+    pub input_len: usize,
+    /// Input length for turns 1+. 0 = fallback to input_len.
+    pub per_turn_input_len: usize,
+    pub output_len: usize,
+    pub seed: u64,
+    pub request_id_prefix: String,
+    pub prefix_sharing_config: Option<PrefixSharingConfig>,
+}
+
+/// Configuration for 3-tier prefix sharing in multi-turn user messages.
+#[derive(Debug, Clone)]
+pub struct PrefixSharingConfig {
+    /// Fraction of per-turn input tokens shared across ALL conversations.
+    pub global_ratio: f64,
+    /// Fraction of per-turn input tokens shared within each conversation.
+    pub conversation_ratio: f64,
+}
+
+/// Generate a deterministic token sequence from allowed tokens using offset+modulo.
+fn make_token_seq(allowed_tokens: &[u32], offset: usize, len: usize) -> Vec<u32> {
+    let at_len = allowed_tokens.len();
+    (0..len).map(|i| allowed_tokens[(offset + i) % at_len]).collect()
+}
+
+/// Generate synthetic multi-turn conversations with random user messages.
+///
+/// Each conversation has `num_turns` turns, each with a random user prompt
+/// of `input_len` tokens and `output_len` expected output tokens.
+pub fn generate_multi_turn_random(
+    tokenizer: &TokenizerKind,
+    cfg: &MultiTurnRandomConfig,
+) -> Result<Vec<MultiTurnConversation>> {
+    let num_conversations = cfg.num_conversations;
+    let min_turns = cfg.min_turns;
+    let max_turns = cfg.max_turns;
+    let prefix_len = cfg.prefix_len;
+    let input_len = cfg.input_len;
+    let output_len = cfg.output_len;
+    let seed = cfg.seed;
+    let request_id_prefix = &cfg.request_id_prefix;
+    let allowed_tokens = tokenizer.get_allowed_tokens();
+    if allowed_tokens.is_empty() {
+        return Err(BenchError::Tokenizer("No allowed tokens found".into()));
+    }
+
+    let vocab_size = tokenizer.vocab_size() as usize;
+    let num_special = tokenizer.num_special_tokens_to_add();
+    let real_input_len = input_len.saturating_sub(num_special);
+    let real_per_turn_len = if cfg.per_turn_input_len > 0 {
+        cfg.per_turn_input_len.saturating_sub(num_special)
+    } else {
+        real_input_len
+    };
+
+    if real_input_len < 1 {
+        return Err(BenchError::Config(format!(
+            "--random-input-len too small: with {num_special} special tokens, \
+             effective input length is {real_input_len}"
+        )));
+    }
+    if real_per_turn_len < 1 {
+        return Err(BenchError::Config(format!(
+            "--per-turn-input-len too small: with {num_special} special tokens, \
+             effective per-turn input length is {real_per_turn_len}"
+        )));
+    }
+
+    // Prefix sharing mode: generate 3-tier prefixed messages
+    let mut rng = StdRng::seed_from_u64(seed);
+    if let Some(ref ps_cfg) = cfg.prefix_sharing_config {
+        return generate_prefix_sharing_conversations(
+            tokenizer,
+            cfg,
+            ps_cfg,
+            &allowed_tokens,
+            &mut rng,
+        );
+    }
+    let shared_prefix_text =
+        generate_shared_prefix_text(tokenizer, &allowed_tokens, prefix_len, seed)?;
+
+    // Pre-generate per-conversation turn counts and per-turn offsets deterministically.
+    // Turn counts are drawn first so the RNG sequence is stable regardless of vocab_size.
+    let conv_turn_counts: Vec<usize> = (0..num_conversations)
+        .map(|_| {
+            if min_turns == max_turns {
+                min_turns
+            } else {
+                rng.random_range(min_turns..=max_turns)
+            }
+        })
+        .collect();
+
+    let offsets: Vec<Vec<usize>> = conv_turn_counts
+        .iter()
+        .map(|&n| (0..n).map(|_| rng.random_range(0..vocab_size)).collect())
+        .collect();
+
+    // Parallel generation across conversations
+    offsets
+        .par_iter()
+        .enumerate()
+        .map(|(conv_idx, conv_offsets)| {
+            let mut turns = Vec::with_capacity(conv_offsets.len());
+            for (turn_idx, &offset) in conv_offsets.iter().enumerate() {
+                let target_len = if turn_idx == 0 {
+                    real_input_len
+                } else {
+                    real_per_turn_len
+                };
+                // Use max_turns stride to keep offsets unique across variable-length convs
+                let inner_seq = make_token_seq(
+                    &allowed_tokens,
+                    offset + conv_idx * max_turns + turn_idx,
+                    target_len,
+                );
+
+                let (prompt, adjusted) =
+                    gen_prompt_to_target_len(tokenizer, &inner_seq, target_len)?;
+                let (prompt, token_len) = if turn_idx == 0 && !shared_prefix_text.is_empty() {
+                    let combined = format!("{}{}", &*shared_prefix_text, prompt);
+                    let token_len = tokenizer.encode(&combined, false)?.len();
+                    (combined, token_len)
+                } else {
+                    (prompt, adjusted.len())
+                };
+
+                turns.push(ConversationTurn {
+                    user_message: Arc::from(prompt),
+                    user_message_len: token_len,
+                    expected_output_len: output_len,
+                });
+            }
+
+            Ok(MultiTurnConversation {
+                conversation_id: format!("{request_id_prefix}conv-{conv_idx}"),
+                turns,
+            })
+        })
+        .collect()
+}
+
+/// Generate conversations with 3-tier prefix sharing.
+///
+/// Each turn's user message = [global_prefix][conversation_prefix][unique_suffix].
+/// No history accumulation — each turn sends only its own fixed-length message.
+fn generate_prefix_sharing_conversations(
+    tokenizer: &TokenizerKind,
+    cfg: &MultiTurnRandomConfig,
+    ps_cfg: &PrefixSharingConfig,
+    allowed_tokens: &[u32],
+    rng: &mut StdRng,
+) -> Result<Vec<MultiTurnConversation>> {
+    let num_conversations = cfg.num_conversations;
+    let min_turns = cfg.min_turns;
+    let max_turns = cfg.max_turns;
+    let prefix_len = cfg.prefix_len;
+    let output_len = cfg.output_len;
+    let request_id_prefix = &cfg.request_id_prefix;
+
+    let num_special = tokenizer.num_special_tokens_to_add();
+    let real_input_len = cfg.input_len.saturating_sub(num_special);
+    let real_per_turn_len = if cfg.per_turn_input_len > 0 {
+        cfg.per_turn_input_len.saturating_sub(num_special)
+    } else {
+        real_input_len
+    };
+
+    // Compute segment lengths from turn-0 (real_input_len) so the shared prefix
+    // bytes stay byte-identical across all turns regardless of per_turn_input_len.
+    let global_len = (real_input_len as f64 * ps_cfg.global_ratio).floor() as usize;
+    let conv_len = (real_input_len as f64 * ps_cfg.conversation_ratio).floor() as usize;
+    let unique_len = real_input_len.saturating_sub(global_len + conv_len);
+
+    // Validate that turns 1+ still have room for a non-empty unique suffix
+    if real_per_turn_len <= global_len + conv_len {
+        return Err(BenchError::Config(format!(
+            "--per-turn-input-len ({real_per_turn_len} after special tokens) is too small: \
+             global_len={global_len} + conv_len={conv_len} already fills the budget. \
+             Increase --per-turn-input-len or reduce prefix ratios."
+        )));
+    }
+
+    let at_len = allowed_tokens.len();
+    let shared_prefix_text =
+        generate_shared_prefix_text(tokenizer, allowed_tokens, prefix_len, cfg.seed)?;
+
+    // Generate global prefix text once
+    let global_text: Arc<str> = if global_len > 0 {
+        let offset: usize = rng.random_range(0..at_len);
+        let seq = make_token_seq(allowed_tokens, offset, global_len);
+        let (text, _) = gen_prompt_to_target_len(tokenizer, &seq, global_len)?;
+        Arc::from(text)
+    } else {
+        Arc::from("")
+    };
+
+    // Generate per-conversation prefix texts
+    let conv_texts: Vec<Arc<str>> = if conv_len > 0 {
+        let mut texts = Vec::with_capacity(num_conversations);
+        for conv_idx in 0..num_conversations {
+            let offset: usize = rng.random_range(0..at_len);
+            let seq = make_token_seq(allowed_tokens, offset + conv_idx, conv_len);
+            let (text, _) = gen_prompt_to_target_len(tokenizer, &seq, conv_len)?;
+            texts.push(Arc::from(text));
+        }
+        texts
+    } else {
+        vec![Arc::from(""); num_conversations]
+    };
+
+    // Pre-generate per-conversation turn counts and unique offsets deterministically.
+    let vocab_size = tokenizer.vocab_size() as usize;
+    let conv_turn_counts: Vec<usize> = (0..num_conversations)
+        .map(|_| {
+            if min_turns == max_turns {
+                min_turns
+            } else {
+                rng.random_range(min_turns..=max_turns)
+            }
+        })
+        .collect();
+
+    let unique_offsets: Vec<Vec<usize>> = conv_turn_counts
+        .iter()
+        .map(|&n| (0..n).map(|_| rng.random_range(0..vocab_size)).collect())
+        .collect();
+
+    // Parallel generation across conversations
+    unique_offsets
+        .par_iter()
+        .enumerate()
+        .map(|(conv_idx, conv_offsets)| {
+            let mut turns = Vec::with_capacity(conv_offsets.len());
+            for (turn_idx, &offset) in conv_offsets.iter().enumerate() {
+                // Turn 0 uses unique_len derived from real_input_len;
+                // turns 1+ use per-turn unique_len (prefix bytes stay identical).
+                let turn_unique_len = if turn_idx == 0 {
+                    unique_len
+                } else {
+                    real_per_turn_len.saturating_sub(global_len + conv_len)
+                };
+
+                // Generate unique suffix
+                let unique_text = if turn_unique_len > 0 {
+                    let seq = make_token_seq(
+                        allowed_tokens,
+                        offset + conv_idx * max_turns + turn_idx,
+                        turn_unique_len,
+                    );
+                    let (text, _) = gen_prompt_to_target_len(tokenizer, &seq, turn_unique_len)?;
+                    text
+                } else {
+                    String::new()
+                };
+
+                // Concatenate: optional random prefix + global + conversation + unique.
+                // Prefix-sharing mode sends each turn independently, so the random
+                // prefix must be included on every turn to be present in every request.
+                let combined = format!(
+                    "{}{}{}{}",
+                    &*shared_prefix_text, &*global_text, &*conv_texts[conv_idx], unique_text
+                );
+                // Re-encode to get actual token count (BPE boundary effects)
+                let token_len = tokenizer.encode(&combined, false)?.len();
+
+                turns.push(ConversationTurn {
+                    user_message: Arc::from(combined),
+                    user_message_len: token_len,
+                    expected_output_len: output_len,
+                });
+            }
+
+            Ok(MultiTurnConversation {
+                conversation_id: format!("{request_id_prefix}conv-{conv_idx}"),
+                turns,
+            })
+        })
+        .collect()
+}
+
+fn generate_shared_prefix_text(
+    tokenizer: &TokenizerKind,
+    allowed_tokens: &[u32],
+    prefix_len: usize,
+    seed: u64,
+) -> Result<Arc<str>> {
+    if prefix_len == 0 {
+        return Ok(Arc::from(""));
+    }
+
+    let mut rng = StdRng::seed_from_u64(seed.wrapping_add(0xDEAD));
+    let tokens: Vec<u32> = (0..prefix_len)
+        .map(|_| allowed_tokens[rng.random_range(0..allowed_tokens.len())])
+        .collect();
+    let (text, _) = gen_prompt_to_target_len(tokenizer, &tokens, prefix_len)?;
+    Ok(Arc::from(text))
+}
+
+/// Load multi-turn conversations from a ShareGPT dataset.
+///
+/// Walks ALL turns in each entry (not just first 2). Filters entries
+/// with at least 4 messages (2 user + 2 assistant = 2 real turns).
+pub fn load_sharegpt_multi_turn(
+    tokenizer: &TokenizerKind,
+    dataset_path: &str,
+    num_conversations: usize,
+    output_len_override: Option<usize>,
+    max_turns: Option<usize>,
+    seed: u64,
+    request_id_prefix: &str,
+) -> Result<Vec<MultiTurnConversation>> {
+    let content = std::fs::read_to_string(dataset_path).map_err(|e| {
+        BenchError::Config(format!(
+            "Failed to read ShareGPT file '{dataset_path}': {e}"
+        ))
+    })?;
+
+    let data: serde_json::Value = serde_json::from_str(&content)
+        .map_err(|e| BenchError::Config(format!("Invalid JSON in ShareGPT file: {e}")))?;
+
+    let entries = data
+        .as_array()
+        .ok_or_else(|| BenchError::Config("ShareGPT file must contain a JSON array".into()))?;
+
+    // Filter entries with at least 4 messages (2 turns: user+assistant+user+assistant)
+    let mut filtered: Vec<&serde_json::Value> = entries
+        .iter()
+        .filter(|entry| {
+            entry
+                .get("conversations")
+                .and_then(|c| c.as_array())
+                .map(|a| a.len() >= 4)
+                .unwrap_or(false)
+        })
+        .collect();
+
+    if filtered.is_empty() {
+        return Err(BenchError::Config(
+            "No valid multi-turn entries in ShareGPT file (need at least 4 messages per entry)"
+                .into(),
+        ));
+    }
+
+    // Shuffle
+    let mut rng = StdRng::seed_from_u64(seed);
+    filtered.shuffle(&mut rng);
+
+    let mut conversations = Vec::new();
+
+    for entry in &filtered {
+        if conversations.len() >= num_conversations {
+            break;
+        }
+
+        let msgs = entry["conversations"].as_array().unwrap();
+        let mut turns = Vec::new();
+
+        // Walk alternating human/gpt pairs, stopping early once max_turns reached
+        // to avoid tokenizing turns that would be discarded by truncate().
+        let mut i = 0;
+        while i + 1 < msgs.len() {
+            if let Some(m) = max_turns
+                && turns.len() >= m
+            {
+                break;
+            }
+            let from = msgs[i].get("from").and_then(|f| f.as_str()).unwrap_or("");
+            let user_text = msgs[i].get("value").and_then(|v| v.as_str()).unwrap_or("");
+            let assistant_text = msgs[i + 1].get("value").and_then(|v| v.as_str()).unwrap_or("");
+
+            // Expect human then gpt
+            if from != "human" || user_text.is_empty() {
+                i += 1;
+                continue;
+            }
+
+            let user_ids = tokenizer.encode(user_text, false)?;
+            let user_len = user_ids.len();
+
+            let expected_output_len = if let Some(override_len) = output_len_override {
+                override_len
+            } else {
+                let assistant_ids = tokenizer.encode(assistant_text, false)?;
+                assistant_ids.len().max(1)
+            };
+
+            turns.push(ConversationTurn {
+                user_message: Arc::from(user_text),
+                user_message_len: user_len,
+                expected_output_len,
+            });
+
+            i += 2;
+        }
+
+        if turns.len() >= 2 {
+            let conv_idx = conversations.len();
+            conversations.push(MultiTurnConversation {
+                conversation_id: format!("{request_id_prefix}conv-{conv_idx}"),
+                turns,
+            });
+        }
+    }
+
+    if conversations.is_empty() {
+        return Err(BenchError::Config(
+            "No valid multi-turn conversations after filtering ShareGPT dataset.".into(),
+        ));
+    }
+
+    // Oversample if needed
+    if conversations.len() < num_conversations {
+        let original_len = conversations.len();
+        let needed = num_conversations - original_len;
+        for i in 0..needed {
+            let mut conv = conversations[rng.random_range(0..original_len)].clone();
+            conv.conversation_id = format!("{request_id_prefix}conv-{}", original_len + i);
+            conversations.push(conv);
+        }
+        println!(
+            "Oversampled multi-turn conversations from {original_len} to {} total.",
+            conversations.len()
+        );
+    }
+
+    Ok(conversations)
+}
+
+/// Ensure decoded-then-encoded prompt length matches the target.
+fn gen_prompt_to_target_len(
+    tokenizer: &TokenizerKind,
+    token_sequence: &[u32],
+    target_len: usize,
+) -> Result<(String, Vec<u32>)> {
+    let max_retry = 20;
+    let mut tokens = token_sequence.to_vec();
+
+    for retry in 0..=max_retry {
+        let prompt = tokenizer.decode(&tokens, true)?;
+        tokens = tokenizer.encode(&prompt, false)?;
+
+        if retry >= max_retry {
+            // BPE tokenizers can oscillate by ±1 on certain boundaries.
+            // For benchmark random content, accept close-enough and truncate/pad.
+            if tokens.len() > target_len {
+                tokens.truncate(target_len);
+            }
+            // If still short by 1-2 tokens, accept as-is — negligible for benchmarks.
+            // Re-decode after truncation to ensure prompt string matches token vector.
+            let prompt = tokenizer.decode(&tokens, true)?;
+            return Ok((prompt, tokens));
+        }
+
+        if tokens.len() == target_len {
+            return Ok((prompt, tokens));
+        } else if tokens.len() < target_len {
+            let allowed = tokenizer.get_allowed_tokens();
+            let needed = target_len - tokens.len();
+            if allowed.is_empty() {
+                let vocab_size = tokenizer.vocab_size() as usize;
+                for j in 0..needed {
+                    tokens.push(((tokens.len() + j) % vocab_size) as u32);
+                }
+            } else {
+                for j in 0..needed {
+                    tokens.push(allowed[(tokens.len() + j) % allowed.len()]);
+                }
+            }
+        } else {
+            tokens.truncate(target_len);
+        }
+    }
+
+    unreachable!()
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+
+    fn common_prefix_bytes(strings: &[&str]) -> usize {
+        if strings.is_empty() {
+            return 0;
+        }
+        let first = strings[0].as_bytes();
+        let mut len = first.len();
+        for s in &strings[1..] {
+            let b = s.as_bytes();
+            len = len.min(b.len());
+            for i in 0..len {
+                if first[i] != b[i] {
+                    len = i;
+                    break;
+                }
+            }
+        }
+        len
+    }
+
+    #[test]
+    #[ignore]
+    fn test_prefix_sharing_structure() {
+        let tok = crate::tokenizer::load_tokenizer("nvidia/Kimi-K2.5-NVFP4", false, None).unwrap();
+
+        let cfg = MultiTurnRandomConfig {
+            num_conversations: 5,
+            min_turns: 3,
+            max_turns: 3,
+            prefix_len: 0,
+            input_len: 1000,
+            per_turn_input_len: 0,
+            output_len: 100,
+            seed: 42,
+            request_id_prefix: "test-".to_string(),
+            prefix_sharing_config: Some(PrefixSharingConfig {
+                global_ratio: 0.1,
+                conversation_ratio: 0.8,
+            }),
+        };
+
+        let conversations = generate_multi_turn_random(&tok, &cfg).unwrap();
+        assert_eq!(conversations.len(), 5);
+
+        let messages: Vec<Vec<&str>> = conversations
+            .iter()
+            .map(|c| c.turns.iter().map(|t| &*t.user_message).collect())
+            .collect();
+
+        // 1. Global prefix: all messages share a common prefix
+        let all_msgs: Vec<&str> = messages.iter().flat_map(|v| v.iter().copied()).collect();
+        let global_prefix = common_prefix_bytes(&all_msgs);
+        println!("Global prefix bytes: {global_prefix}");
+        assert!(global_prefix > 0, "Global prefix must be non-empty");
+
+        // 2. Conversation prefix: turns within same conversation share more
+        for (i, conv_msgs) in messages.iter().enumerate() {
+            let conv_prefix = common_prefix_bytes(conv_msgs);
+            println!("Conv {i} prefix bytes: {conv_prefix} (global: {global_prefix})");
+            assert!(
+                conv_prefix > global_prefix,
+                "Conv prefix ({conv_prefix}) must exceed global prefix ({global_prefix})"
+            );
+        }
+
+        // 3. Different conversations diverge after global prefix
+        let cross = common_prefix_bytes(&[messages[0][0], messages[1][0]]);
+        let within = common_prefix_bytes(&messages[0]);
+        println!("Cross-conv prefix: {cross}, within-conv prefix: {within}");
+        assert!(
+            cross < within,
+            "Cross-conv ({cross}) must be < within-conv ({within})"
+        );
+
+        // 4. Turns within same conversation are not identical (unique suffix)
+        for (i, conv_msgs) in messages.iter().enumerate() {
+            for a in 0..conv_msgs.len() {
+                for b in (a + 1)..conv_msgs.len() {
+                    assert_ne!(
+                        conv_msgs[a], conv_msgs[b],
+                        "Conv {i} turn {a} and {b} must differ"
+                    );
+                }
+            }
+        }
+
+        // 5. Token lengths approximately match target
+        for (i, conv) in conversations.iter().enumerate() {
+            for (j, turn) in conv.turns.iter().enumerate() {
+                let diff = (turn.user_message_len as i64 - 1000).abs();
+                println!(
+                    "Conv {i} turn {j}: {} tokens (diff {diff})",
+                    turn.user_message_len
+                );
+                assert!(
+                    diff <= 10,
+                    "Token len {} too far from 1000",
+                    turn.user_message_len
+                );
+            }
+        }
+
+        println!("All prefix sharing checks passed!");
+    }
+
+    #[test]
+    #[ignore]
+    fn test_per_turn_input_len_default_mode() {
+        let tok = crate::tokenizer::load_tokenizer("nvidia/Kimi-K2.5-NVFP4", false, None).unwrap();
+
+        let cfg = MultiTurnRandomConfig {
+            num_conversations: 4,
+            min_turns: 3,
+            max_turns: 3,
+            prefix_len: 0,
+            input_len: 512,
+            per_turn_input_len: 128,
+            output_len: 64,
+            seed: 1,
+            request_id_prefix: "test-".to_string(),
+            prefix_sharing_config: None,
+        };
+
+        let conversations = generate_multi_turn_random(&tok, &cfg).unwrap();
+        assert_eq!(conversations.len(), 4);
+
+        for (i, conv) in conversations.iter().enumerate() {
+            assert_eq!(conv.turns.len(), 3);
+            for (j, turn) in conv.turns.iter().enumerate() {
+                let expected = if j == 0 { 512usize } else { 128usize };
+                let diff = (turn.user_message_len as i64 - expected as i64).abs();
+                println!(
+                    "Conv {i} turn {j}: {} tokens (expected ~{expected}, diff {diff})",
+                    turn.user_message_len
+                );
+                assert!(
+                    diff <= 5,
+                    "Conv {i} turn {j}: token len {} too far from {expected}",
+                    turn.user_message_len
+                );
+            }
+        }
+        println!("per_turn_input_len default-mode checks passed!");
+    }
+
+    #[test]
+    #[ignore]
+    fn test_variable_turns_range() {
+        let tok = crate::tokenizer::load_tokenizer("nvidia/Kimi-K2.5-NVFP4", false, None).unwrap();
+
+        let cfg = MultiTurnRandomConfig {
+            num_conversations: 50,
+            min_turns: 2,
+            max_turns: 5,
+            prefix_len: 0,
+            input_len: 256,
+            per_turn_input_len: 0,
+            output_len: 32,
+            seed: 7,
+            request_id_prefix: "test-".to_string(),
+            prefix_sharing_config: None,
+        };
+
+        let conversations = generate_multi_turn_random(&tok, &cfg).unwrap();
+        assert_eq!(conversations.len(), 50);
+
+        let mut distinct_counts = std::collections::HashSet::new();
+        for conv in &conversations {
+            let n = conv.turns.len();
+            assert!((2..=5).contains(&n), "turn count {n} out of [2,5]");
+            distinct_counts.insert(n);
+        }
+        assert!(
+            distinct_counts.len() >= 2,
+            "expected at least 2 distinct turn counts, got {distinct_counts:?}"
+        );
+        println!("variable_turns_range checks passed! counts: {distinct_counts:?}");
+    }
+
+    #[test]
+    #[ignore]
+    fn test_variable_turns_fixed() {
+        let tok = crate::tokenizer::load_tokenizer("nvidia/Kimi-K2.5-NVFP4", false, None).unwrap();
+
+        let cfg = MultiTurnRandomConfig {
+            num_conversations: 10,
+            min_turns: 4,
+            max_turns: 4,
+            prefix_len: 0,
+            input_len: 256,
+            per_turn_input_len: 0,
+            output_len: 32,
+            seed: 42,
+            request_id_prefix: "test-".to_string(),
+            prefix_sharing_config: None,
+        };
+
+        let conversations = generate_multi_turn_random(&tok, &cfg).unwrap();
+        for conv in &conversations {
+            assert_eq!(conv.turns.len(), 4, "expected exactly 4 turns");
+        }
+        println!("variable_turns_fixed checks passed!");
+    }
+
+    #[test]
+    #[ignore]
+    fn test_per_turn_input_len_prefix_sharing() {
+        let tok = crate::tokenizer::load_tokenizer("nvidia/Kimi-K2.5-NVFP4", false, None).unwrap();
+
+        // Turn 0 input_len=1000, turns 1+ per_turn_input_len=600
+        // global_len ≈ 100 (10%), conv_len ≈ 800 (80%), unique ≈ 100
+        // per-turn unique ≈ 600 - 900 = negative → would error; use smaller ratios
+        // global=0.05 (50), conv=0.5 (500), unique_t0=450, unique_t1=600-550=50
+        let cfg = MultiTurnRandomConfig {
+            num_conversations: 4,
+            min_turns: 3,
+            max_turns: 3,
+            prefix_len: 0,
+            input_len: 1000,
+            per_turn_input_len: 600,
+            output_len: 64,
+            seed: 3,
+            request_id_prefix: "test-".to_string(),
+            prefix_sharing_config: Some(PrefixSharingConfig {
+                global_ratio: 0.05,
+                conversation_ratio: 0.50,
+            }),
+        };
+
+        let conversations = generate_multi_turn_random(&tok, &cfg).unwrap();
+        assert_eq!(conversations.len(), 4);
+
+        let messages: Vec<Vec<&str>> = conversations
+            .iter()
+            .map(|c| c.turns.iter().map(|t| &*t.user_message).collect())
+            .collect();
+
+        // Global prefix bytes shared across all turns of all conversations
+        let all_msgs: Vec<&str> = messages.iter().flat_map(|v| v.iter().copied()).collect();
+        let global_prefix = common_prefix_bytes(&all_msgs);
+        assert!(global_prefix > 0, "Global prefix must be non-empty");
+
+        // Within each conversation, prefix grows (conv prefix longer than global)
+        for (i, conv_msgs) in messages.iter().enumerate() {
+            let conv_prefix = common_prefix_bytes(conv_msgs);
+            assert!(
+                conv_prefix > global_prefix,
+                "Conv {i}: conv_prefix ({conv_prefix}) must exceed global ({global_prefix})"
+            );
+        }
+
+        // Turn 0 length ≈ 1000, turns 1+ ≈ 600
+        for (i, conv) in conversations.iter().enumerate() {
+            for (j, turn) in conv.turns.iter().enumerate() {
+                let expected = if j == 0 { 1000usize } else { 600usize };
+                let diff = (turn.user_message_len as i64 - expected as i64).abs();
+                println!(
+                    "Conv {i} turn {j}: {} tokens (expected ~{expected})",
+                    turn.user_message_len
+                );
+                assert!(
+                    diff <= 10,
+                    "Conv {i} turn {j}: token len {} too far from {expected}",
+                    turn.user_message_len
+                );
+            }
+        }
+        println!("per_turn_input_len prefix-sharing checks passed!");
+    }
+}
@@ -0,0 +1,148 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+//! Prefix repetition dataset: N distinct shared prefixes, each reused by
+//! `num_prompts / num_prefixes` requests with a fresh random suffix.
+//! The standard prefix-cache stress workload; mirrors Python's
+//! `PrefixRepetitionRandomDataset`.
+
+use std::sync::Arc;
+
+use rand::rngs::StdRng;
+use rand::seq::SliceRandom;
+use rand::{Rng, SeedableRng};
+use rayon::prelude::*;
+
+use super::SampleRequest;
+use super::random::gen_prompt_decode_to_target_len;
+use crate::error::{BenchError, Result};
+use crate::tokenizer::TokenizerKind;
+
+/// Generate the prefix repetition dataset.
+///
+/// Like Python, `num_requests % num_prefixes` remainder requests are dropped:
+/// the total is `(num_requests / num_prefixes) * num_prefixes`.
+pub fn generate_prefix_repetition_dataset(
+    tokenizer: &TokenizerKind,
+    num_requests: usize,
+    prefix_len: usize,
+    suffix_len: usize,
+    num_prefixes: usize,
+    output_len: usize,
+    seed: u64,
+    request_id_prefix: &str,
+    disable_shuffle: bool,
+) -> Result<Vec<SampleRequest>> {
+    let prompts_per_prefix = num_requests / num_prefixes;
+    if prompts_per_prefix == 0 {
+        return Err(BenchError::Config(format!(
+            "num_prompts ({num_requests}) must be >= num_prefixes ({num_prefixes})"
+        )));
+    }
+    let total = prompts_per_prefix * num_prefixes;
+    if total != num_requests {
+        println!(
+            "prefix_repetition: generating {total} requests \
+             ({num_prefixes} prefixes x {prompts_per_prefix} prompts each; \
+             {} dropped to divide evenly)",
+            num_requests - total
+        );
+    }
+
+    let allowed_tokens = tokenizer.get_allowed_tokens();
+    if allowed_tokens.is_empty() {
+        return Err(BenchError::Tokenizer("No allowed tokens found".into()));
+    }
+    let allowed_ref = &allowed_tokens;
+
+    // Exact-length random token block: decode -> re-encode -> converge to target.
+    let gen_block = |target_len: usize, item_seed: u64| -> Result<Vec<u32>> {
+        let mut rng = StdRng::seed_from_u64(item_seed);
+        let tokens: Vec<u32> = (0..target_len)
+            .map(|_| allowed_ref[rng.random_range(0..allowed_ref.len())])
+            .collect();
+        let (_, adjusted) =
+            gen_prompt_decode_to_target_len(tokenizer, &tokens, target_len, false, allowed_ref)?;
+        Ok(adjusted)
+    };
+
+    // Generate the shared prefixes (one per group), then suffixes in parallel.
+    let prefixes: Vec<Vec<u32>> = (0..num_prefixes)
+        .map(|p| gen_block(prefix_len, seed.wrapping_add(0xF1F0).wrapping_add(p as u64)))
+        .collect::<Result<Vec<_>>>()?;
+
+    let rid_prefix = request_id_prefix.to_string();
+    let mut requests: Vec<SampleRequest> = (0..total)
+        .into_par_iter()
+        .map(|i| {
+            let prefix_tokens = &prefixes[i / prompts_per_prefix];
+            let suffix_tokens = gen_block(suffix_len, seed.wrapping_add(0xBEEF + i as u64))?;
+
+            let mut combined = Vec::with_capacity(prefix_tokens.len() + suffix_tokens.len());
+            combined.extend_from_slice(prefix_tokens);
+            combined.extend_from_slice(&suffix_tokens);
+            let prompt = tokenizer.decode(&combined, true)?;
+
+            Ok(SampleRequest {
+                prompt: Arc::from(prompt),
+                prompt_len: combined.len(),
+                expected_output_len: output_len,
+                request_id: Some(format!("{rid_prefix}{i}")),
+                ..Default::default()
+            })
+        })
+        .collect::<Result<Vec<_>>>()?;
+
+    // Interleave prefixes (Python shuffles too) so one prefix group isn't sent
+    // as a contiguous burst.
+    if !disable_shuffle {
+        let mut rng = StdRng::seed_from_u64(seed);
+        requests.shuffle(&mut rng);
+    }
+
+    Ok(requests)
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+
+    /// gpt2 via built-in tiktoken encoding — loads without network access.
+    fn test_tokenizer() -> TokenizerKind {
+        crate::tokenizer::load_tokenizer("gpt2", false, None)
+            .expect("gpt2 built-in tiktoken should always load without network")
+    }
+
+    #[test]
+    fn test_prefix_repetition_structure() {
+        let tok = test_tokenizer();
+        // 7 requests / 3 prefixes -> 2 per prefix, 6 total (remainder dropped like Python)
+        let reqs = generate_prefix_repetition_dataset(&tok, 7, 32, 16, 3, 64, 0, "t-", true)
+            .expect("generation should succeed");
+        assert_eq!(reqs.len(), 6);
+        assert!(reqs.iter().all(|r| r.expected_output_len == 64));
+        // Exact-length blocks: prompt_len == prefix + suffix
+        assert!(
+            reqs.iter().all(|r| r.prompt_len == 32 + 16),
+            "lens: {:?}",
+            reqs.iter().map(|r| r.prompt_len).collect::<Vec<_>>()
+        );
+        // Consecutive pairs (shuffle disabled) share a common prefix; requests
+        // from different groups don't.
+        let common = |a: &str, b: &str| -> usize {
+            a.bytes().zip(b.bytes()).take_while(|(x, y)| x == y).count()
+        };
+        let same_group = common(&reqs[0].prompt, &reqs[1].prompt);
+        let diff_group = common(&reqs[0].prompt, &reqs[2].prompt);
+        assert!(
+            same_group > diff_group,
+            "same-group shared prefix ({same_group}) should exceed cross-group ({diff_group})"
+        );
+    }
+
+    #[test]
+    fn test_prefix_repetition_too_few_requests_errors() {
+        let tok = test_tokenizer();
+        assert!(generate_prefix_repetition_dataset(&tok, 2, 32, 16, 3, 64, 0, "t-", true).is_err());
+    }
+}
@@ -0,0 +1,491 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+use std::sync::Arc;
+
+use rayon::prelude::*;
+
+use super::SampleRequest;
+use crate::config::RangeRatio;
+use crate::error::{BenchError, Result};
+use crate::tokenizer::TokenizerKind;
+
+/// Generate random dataset with rayon parallelism.
+///
+/// This is the key performance win — Python does sequential tokenizer calls
+/// while Rust parallelizes across CPU cores with native tokenizer speed.
+///
+/// Mirrors Python's RandomDataset.sample() from datasets.py:470-560.
+pub fn generate_random_dataset(
+    tokenizer: &TokenizerKind,
+    num_requests: usize,
+    input_len: usize,
+    output_len: usize,
+    prefix_len: usize,
+    range_ratio: RangeRatio,
+    cache_hit_fraction: f64,
+    cache_ratio: f64,
+    seed: u64,
+    request_id_prefix: &str,
+    use_token_ids: bool,
+    batch_size: usize,
+) -> Result<Vec<SampleRequest>> {
+    let vocab_size = tokenizer.vocab_size();
+    let allowed_tokens = tokenizer.get_allowed_tokens();
+    if allowed_tokens.is_empty() {
+        return Err(BenchError::Tokenizer("No allowed tokens found".into()));
+    }
+
+    if batch_size > 1 && use_token_ids {
+        return Err(BenchError::Config(
+            "--random-batch-size > 1 is not supported with --prompt-token-ids".into(),
+        ));
+    }
+
+    let num_special = tokenizer.num_special_tokens_to_add();
+    let real_input_len = input_len.saturating_sub(num_special);
+
+    // Python semantics: sample uniformly from [len*(1-r), len*(1+r)].
+    let (input_low, input_high) = range_ratio.input_bounds(real_input_len);
+    let (output_low, output_high) = range_ratio.output_bounds(output_len);
+    if !range_ratio.is_fixed() {
+        println!(
+            "Sampling input_len from [{input_low}, {input_high}] and \
+             output_len from [{output_low}, {output_high}]"
+        );
+    }
+
+    // Bimodal prefix-cache mode: a fraction of prompts (warm) reuse a shared cached
+    // prefix covering `cache_ratio` of their length; the rest (cold) are fully unique.
+    // Models e.g. "80% of prompts have 95% of input cached" with
+    // --random-cache-hit-fraction 0.8 --random-cache-ratio 0.95. In this mode
+    // --random-input-len is the TOTAL prompt length L (the cached prefix is part of L),
+    // and --random-prefix-len is ignored.
+    let bimodal = cache_hit_fraction > 0.0 && cache_ratio > 0.0;
+    if bimodal {
+        if !use_token_ids {
+            return Err(BenchError::Config(
+                "bimodal prefix-cache (--random-cache-hit-fraction) requires --prompt-token-ids \
+                 so warm prompts send identical token IDs and actually hit the prefix cache"
+                    .into(),
+            ));
+        }
+        if cache_hit_fraction > 1.0 || cache_ratio > 1.0 {
+            return Err(BenchError::Config(
+                "--random-cache-hit-fraction and --random-cache-ratio must be in [0, 1]".into(),
+            ));
+        }
+    }
+
+    // Length of the shared cached base prefix.
+    let base_len = if bimodal {
+        ((input_high as f64) * cache_ratio).ceil() as usize
+    } else {
+        prefix_len
+    };
+
+    // Validate (non-bimodal keeps the original check)
+    if !bimodal {
+        let min_total = prefix_len + input_low;
+        if min_total < 1 {
+            return Err(BenchError::Config(format!(
+                "--random-input-len too small: with {num_special} special tokens and \
+                 range_ratio={:?}, minimum total input is {min_total}",
+                range_ratio
+            )));
+        }
+    }
+
+    // Generate the shared base prefix once (sequential, only happens once).
+    let prefix_token_ids = if base_len > 0 {
+        generate_prefix(tokenizer, &allowed_tokens, base_len, seed)?
+    } else {
+        Vec::new()
+    };
+
+    // Pre-generate per-request sampling params using deterministic RNG
+    use rand::rngs::StdRng;
+    use rand::{Rng, SeedableRng};
+    let mut rng = StdRng::seed_from_u64(seed);
+
+    struct RequestParams {
+        cached_len: usize, // tokens taken from the shared base (cache-hittable)
+        suffix_len: usize, // unique tokens appended after the cached prefix
+        output_len: usize,
+        offset: usize,
+    }
+
+    let params: Vec<RequestParams> = (0..num_requests)
+        .map(|_| {
+            let ol = if output_low == output_high {
+                output_low
+            } else {
+                rng.random_range(output_low..=output_high)
+            };
+            let off = rng.random_range(0..vocab_size as usize);
+            if bimodal {
+                // Total length L from the input distribution; prefix is part of L.
+                let l = if input_low == input_high {
+                    input_low
+                } else {
+                    rng.random_range(input_low..=input_high)
+                };
+                let warm = rng.random::<f64>() < cache_hit_fraction;
+                let cached = if warm {
+                    (((l as f64) * cache_ratio).round() as usize).min(base_len).min(l)
+                } else {
+                    0
+                };
+                RequestParams {
+                    cached_len: cached,
+                    suffix_len: l - cached,
+                    output_len: ol,
+                    offset: off,
+                }
+            } else {
+                // Original behavior: full shared prefix + variable unique input.
+                let il = if input_low == input_high {
+                    input_low
+                } else {
+                    rng.random_range(input_low..=input_high)
+                };
+                RequestParams {
+                    cached_len: prefix_len,
+                    suffix_len: il,
+                    output_len: ol,
+                    offset: off,
+                }
+            }
+        })
+        .collect();
+
+    // Phase 1: Generate all token sequences (parallel, fast — just array ops)
+    let prefix_ref = &prefix_token_ids;
+    let allowed_ref = &allowed_tokens;
+    let rid_prefix = request_id_prefix.to_string();
+
+    let token_sequences: Vec<Vec<u32>> = params
+        .par_iter()
+        .enumerate()
+        .map(|(i, p)| {
+            let at_len = allowed_ref.len();
+            let mut seq = Vec::with_capacity(p.cached_len + p.suffix_len);
+            seq.extend_from_slice(&prefix_ref[..p.cached_len]);
+            for j in 0..p.suffix_len {
+                seq.push(allowed_ref[(p.offset + i + j) % at_len]);
+            }
+            seq
+        })
+        .collect();
+
+    let target_lens: Vec<usize> = params.iter().map(|p| p.cached_len + p.suffix_len).collect();
+
+    if use_token_ids {
+        // Fast path: store token IDs directly. The completions backend sends
+        // them as `"prompt": [id1, id2, ...]`, bypassing both client-side decode
+        // and server-side tokenization. Token counts are exact by construction.
+        let result: Vec<SampleRequest> = token_sequences
+            .into_par_iter()
+            .enumerate()
+            .map(|(i, tokens)| SampleRequest {
+                prompt: Arc::from(""),
+                prompt_len: target_lens[i],
+                expected_output_len: params[i].output_len,
+                request_id: Some(format!("{rid_prefix}{i}")),
+                prompt_token_ids: Some(Arc::from(tokens)),
+                ..Default::default()
+            })
+            .collect();
+        Ok(result)
+    } else {
+        // Default path: decode tokens to text, re-encode,
+        // truncate to target length, decode again. Sends text prompts for maximum
+        let result: Vec<SampleRequest> = token_sequences
+            .into_par_iter()
+            .enumerate()
+            .map(|(i, tokens)| {
+                let target = target_lens[i];
+                // decode → encode → truncate → decode
+                let prompt_text = tokenizer.decode(&tokens, true)?;
+                let mut re_encoded = tokenizer.encode(&prompt_text, false)?;
+                re_encoded.truncate(target);
+                let prompt = tokenizer.decode(&re_encoded, true)?;
+                let prompt_len = re_encoded.len();
+                Ok(SampleRequest {
+                    prompt: Arc::from(prompt),
+                    prompt_len,
+                    expected_output_len: params[i].output_len,
+                    request_id: Some(format!("{rid_prefix}{i}")),
+                    ..Default::default()
+                })
+            })
+            .collect::<Result<Vec<_>>>()?;
+        Ok(super::batch_requests(result, batch_size, request_id_prefix))
+    }
+}
+
+fn generate_prefix(
+    tokenizer: &TokenizerKind,
+    allowed_tokens: &[u32],
+    prefix_len: usize,
+    seed: u64,
+) -> Result<Vec<u32>> {
+    use rand::rngs::StdRng;
+    use rand::{Rng, SeedableRng};
+
+    let mut rng = StdRng::seed_from_u64(seed.wrapping_add(0xDEAD));
+    let tokens: Vec<u32> = (0..prefix_len)
+        .map(|_| allowed_tokens[rng.random_range(0..allowed_tokens.len())])
+        .collect();
+
+    let (_, adjusted) =
+        gen_prompt_decode_to_target_len(tokenizer, &tokens, prefix_len, false, allowed_tokens)?;
+    Ok(adjusted)
+}
+
+/// Ensure decoded-then-encoded prompt length matches the target.
+///
+/// Mirrors Python's `gen_prompt_decode_to_target_len` from datasets.py:381-435.
+pub(crate) fn gen_prompt_decode_to_target_len(
+    tokenizer: &TokenizerKind,
+    token_sequence: &[u32],
+    target_len: usize,
+    add_special_tokens: bool,
+    allowed_tokens: &[u32],
+) -> Result<(String, Vec<u32>)> {
+    let max_retry = 20;
+    let mut tokens = token_sequence.to_vec();
+
+    for retry in 0..=max_retry {
+        let prompt = tokenizer.decode(&tokens, true)?;
+        tokens = tokenizer.encode(&prompt, add_special_tokens)?;
+
+        if retry >= max_retry {
+            if tokens.len() != target_len {
+                return Err(BenchError::Tokenizer(format!(
+                    "Token length mismatch after {max_retry} retries: \
+                     target={target_len}, actual={}. \
+                     encode/decode roundtrip cannot converge.",
+                    tokens.len()
+                )));
+            }
+            return Ok((prompt, tokens));
+        }
+
+        if tokens.len() == target_len {
+            return Ok((prompt, tokens));
+        } else if tokens.len() < target_len {
+            // Pad with tokens from the allowed set (UTF-8-safe for tiktoken)
+            let needed = target_len - tokens.len();
+            if allowed_tokens.is_empty() {
+                let vocab_size = tokenizer.vocab_size() as usize;
+                for j in 0..needed {
+                    tokens.push(((tokens.len() + j) % vocab_size) as u32);
+                }
+            } else {
+                for j in 0..needed {
+                    tokens.push(allowed_tokens[(tokens.len() + j) % allowed_tokens.len()]);
+                }
+            }
+        } else {
+            // Truncate
+            tokens.truncate(target_len);
+        }
+    }
+
+    unreachable!()
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+    use crate::tokenizer;
+
+    // Integration test requires a tokenizer, so only run with --ignored
+    #[test]
+    #[ignore]
+    fn test_generate_random_dataset_token_ids() {
+        let tokenizer = tokenizer::load_tokenizer("gpt2", false, None).unwrap();
+        let requests = generate_random_dataset(
+            &tokenizer,
+            10,  // num_requests
+            128, // input_len
+            32,  // output_len
+            0,   // prefix_len
+            RangeRatio {
+                input: 0.0,
+                output: 0.0,
+            }, // range_ratio (0.0 = fixed length)
+            0.0, // cache_hit_fraction (0 = bimodal off)
+            0.0, // cache_ratio
+            42,  // seed
+            "test-",
+            true, // use_token_ids
+            1,    // batch_size
+        )
+        .unwrap();
+
+        assert_eq!(requests.len(), 10);
+        for req in &requests {
+            assert!(req.prompt_token_ids.is_some());
+            assert_eq!(req.prompt_token_ids.as_ref().unwrap().len(), req.prompt_len);
+            assert!(req.prompt_len > 0);
+            assert_eq!(req.expected_output_len, 32);
+        }
+    }
+
+    #[test]
+    #[ignore]
+    fn test_generate_random_dataset_text() {
+        let tokenizer = tokenizer::load_tokenizer("gpt2", false, None).unwrap();
+        let requests = generate_random_dataset(
+            &tokenizer,
+            10,  // num_requests
+            128, // input_len
+            32,  // output_len
+            0,   // prefix_len
+            RangeRatio {
+                input: 0.0,
+                output: 0.0,
+            }, // range_ratio (0.0 = fixed length)
+            0.0, // cache_hit_fraction (0 = bimodal off)
+            0.0, // cache_ratio
+            42,  // seed
+            "test-",
+            false, // use_token_ids = false → text prompts
+            1,     // batch_size
+        )
+        .unwrap();
+
+        assert_eq!(requests.len(), 10);
+        for req in &requests {
+            assert!(req.prompt_token_ids.is_none());
+            assert!(!req.prompt.is_empty());
+            assert!(req.prompt_len > 0);
+            assert!(req.prompt_len <= 128);
+            assert_eq!(req.expected_output_len, 32);
+        }
+    }
+
+    /// Test that generated prompts have EXACT target token length (token ID mode).
+    #[test]
+    #[ignore]
+    fn test_token_length_exact_local() {
+        let tokenizer = tokenizer::load_tokenizer("gpt2", false, None).unwrap();
+        let target_len = 512;
+        let requests = generate_random_dataset(
+            &tokenizer,
+            50,
+            target_len,
+            64,
+            0,
+            RangeRatio {
+                input: 0.0,
+                output: 0.0,
+            },
+            0.0,
+            0.0,
+            123,
+            "len-test-",
+            true,
+            1,
+        )
+        .unwrap();
+
+        for (i, req) in requests.iter().enumerate() {
+            let token_ids = req.prompt_token_ids.as_ref().expect("should have token IDs");
+            assert_eq!(
+                token_ids.len(),
+                target_len,
+                "Request {i}: expected {target_len} token IDs, got {}",
+                token_ids.len()
+            );
+            assert_eq!(req.prompt_len, target_len);
+        }
+    }
+
+    /// Test that tiktoken tokenizer produces exact target token lengths (token ID mode).
+    #[test]
+    #[ignore]
+    fn test_token_length_exact_tiktoken() {
+        // Use Qwen2.5 which has a tiktoken-format tokenizer
+        let tokenizer = tokenizer::load_tokenizer("Qwen/Qwen2.5-0.5B", false, None);
+        let tokenizer = match tokenizer {
+            Ok(t) => t,
+            Err(e) => {
+                eprintln!("Skipping tiktoken test (tokenizer unavailable): {e}");
+                return;
+            }
+        };
+
+        // Verify it's actually a tiktoken tokenizer or local — either way test convergence
+        let target_len = 256;
+        let requests = generate_random_dataset(
+            &tokenizer,
+            20,
+            target_len,
+            32,
+            0,
+            RangeRatio {
+                input: 0.0,
+                output: 0.0,
+            },
+            0.0,
+            0.0,
+            42,
+            "tiktoken-test-",
+            true,
+            1,
+        )
+        .unwrap();
+
+        for (i, req) in requests.iter().enumerate() {
+            let token_ids = req.prompt_token_ids.as_ref().expect("should have token IDs");
+            assert_eq!(
+                token_ids.len(),
+                target_len,
+                "Request {i}: expected {target_len} token IDs, got {}",
+                token_ids.len()
+            );
+            assert_eq!(req.prompt_len, target_len);
+        }
+    }
+
+    /// Test encode/decode roundtrip stability for tiktoken.
+    /// After one decode→encode cycle with UTF-8-safe tokens, length must not drift.
+    #[test]
+    #[ignore]
+    fn test_tiktoken_roundtrip_stability() {
+        let tokenizer = tokenizer::load_tokenizer("Qwen/Qwen2.5-0.5B", false, None);
+        let tokenizer = match tokenizer {
+            Ok(t) => t,
+            Err(e) => {
+                eprintln!("Skipping roundtrip test (tokenizer unavailable): {e}");
+                return;
+            }
+        };
+
+        let allowed = tokenizer.get_allowed_tokens();
+        assert!(!allowed.is_empty(), "allowed tokens should not be empty");
+
+        // Build a sequence from allowed tokens only
+        use rand::rngs::StdRng;
+        use rand::{Rng, SeedableRng};
+        let mut rng = StdRng::seed_from_u64(99);
+        let seq: Vec<u32> = (0..512).map(|_| allowed[rng.random_range(0..allowed.len())]).collect();
+
+        let decoded = tokenizer.decode(&seq, true).unwrap();
+        let re_encoded = tokenizer.encode(&decoded, false).unwrap();
+        let re_decoded = tokenizer.decode(&re_encoded, true).unwrap();
+        let re_re_encoded = tokenizer.encode(&re_decoded, false).unwrap();
+
+        // After first cycle, length should stabilize
+        assert_eq!(
+            re_encoded.len(),
+            re_re_encoded.len(),
+            "Roundtrip should stabilize: first re-encode={}, second re-encode={}",
+            re_encoded.len(),
+            re_re_encoded.len()
+        );
+    }
+}
@@ -0,0 +1,657 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+use std::io::Cursor;
+use std::sync::Arc;
+
+use base64::Engine as _;
+use rand::rngs::StdRng;
+use rand::{Rng, SeedableRng};
+use rayon::prelude::*;
+
+use super::SampleRequest;
+use crate::error::{BenchError, Result};
+use crate::tokenizer::TokenizerKind;
+
+/// A bucket key: (height, width, num_frames). num_frames=1 means image, >1 means video.
+#[derive(Debug, Clone)]
+pub struct MmBucketKey {
+    pub height: u32,
+    pub width: u32,
+    pub num_frames: u32,
+}
+
+/// Per-modality hard caps.
+#[derive(Debug, Clone)]
+pub struct MmLimitPerPrompt {
+    pub image: usize,
+    pub video: usize,
+}
+
+impl Default for MmLimitPerPrompt {
+    fn default() -> Self {
+        Self {
+            image: 255,
+            video: 1,
+        }
+    }
+}
+
+/// Parse the limit-mm-per-prompt JSON string, e.g. `{"image": 3, "video": 0}`.
+pub fn parse_limit_mm_per_prompt(s: &str) -> Result<MmLimitPerPrompt> {
+    let v: serde_json::Value = serde_json::from_str(s)
+        .map_err(|e| BenchError::Config(format!("Invalid --random-mm-limit-mm-per-prompt: {e}")))?;
+    let obj = v.as_object().ok_or_else(|| {
+        BenchError::Config("--random-mm-limit-mm-per-prompt must be a JSON object".into())
+    })?;
+    let image = obj.get("image").and_then(|v| v.as_u64()).unwrap_or(255) as usize;
+    let video = obj.get("video").and_then(|v| v.as_u64()).unwrap_or(1) as usize;
+    Ok(MmLimitPerPrompt { image, video })
+}
+
+/// Parse the bucket config string in Python-style syntax.
+///
+/// Accepts: `{(256,256,1): 0.5, (720,1280,1): 0.5}`
+/// Each key is `(height, width, num_frames)` and value is the probability weight.
+pub fn parse_bucket_config(s: &str) -> Result<Vec<(MmBucketKey, f64)>> {
+    let trimmed = s.trim();
+    let inner = trimmed
+        .strip_prefix('{')
+        .and_then(|s| s.strip_suffix('}'))
+        .ok_or_else(|| BenchError::Config("Bucket config must be wrapped in {}".into()))?;
+
+    let mut buckets = Vec::new();
+    let chars: Vec<char> = inner.chars().collect();
+    let len = chars.len();
+    let mut i = 0;
+
+    while i < len {
+        // Skip whitespace and commas
+        while i < len && (chars[i].is_whitespace() || chars[i] == ',') {
+            i += 1;
+        }
+        if i >= len {
+            break;
+        }
+
+        // Expect '('
+        if chars[i] != '(' {
+            return Err(BenchError::Config(format!(
+                "Expected '(' in bucket config at position {i}"
+            )));
+        }
+        i += 1;
+
+        // Read until ')'
+        let tuple_start = i;
+        while i < len && chars[i] != ')' {
+            i += 1;
+        }
+        if i >= len {
+            return Err(BenchError::Config("Unclosed '(' in bucket config".into()));
+        }
+        let tuple_str: String = chars[tuple_start..i].iter().collect();
+        i += 1; // skip ')'
+
+        // Skip whitespace, then expect ':'
+        while i < len && chars[i].is_whitespace() {
+            i += 1;
+        }
+        if i >= len || chars[i] != ':' {
+            return Err(BenchError::Config(
+                "Expected ':' after tuple in bucket config".into(),
+            ));
+        }
+        i += 1;
+
+        // Skip whitespace
+        while i < len && chars[i].is_whitespace() {
+            i += 1;
+        }
+
+        // Read the probability value until ',' or end
+        let val_start = i;
+        while i < len && chars[i] != ',' {
+            i += 1;
+        }
+        let val_str: String = chars[val_start..i].iter().collect();
+
+        // Parse tuple
+        let parts: Vec<&str> = tuple_str.split(',').collect();
+        if parts.len() != 3 {
+            return Err(BenchError::Config(format!(
+                "Bucket key must have 3 values (height,width,num_frames), got: ({tuple_str})"
+            )));
+        }
+
+        let height: u32 = parts[0].trim().parse().map_err(|_| {
+            BenchError::Config(format!(
+                "Invalid height in bucket config: '{}'",
+                parts[0].trim()
+            ))
+        })?;
+        let width: u32 = parts[1].trim().parse().map_err(|_| {
+            BenchError::Config(format!(
+                "Invalid width in bucket config: '{}'",
+                parts[1].trim()
+            ))
+        })?;
+        let num_frames: u32 = parts[2].trim().parse().map_err(|_| {
+            BenchError::Config(format!(
+                "Invalid num_frames in bucket config: '{}'",
+                parts[2].trim()
+            ))
+        })?;
+        let prob: f64 = val_str.trim().parse().map_err(|_| {
+            BenchError::Config(format!(
+                "Invalid probability in bucket config: '{}'",
+                val_str.trim()
+            ))
+        })?;
+
+        if prob < 0.0 {
+            return Err(BenchError::Config(format!(
+                "Bucket probability must be non-negative, got: {prob}"
+            )));
+        }
+
+        buckets.push((
+            MmBucketKey {
+                height,
+                width,
+                num_frames,
+            },
+            prob,
+        ));
+    }
+
+    if buckets.is_empty() {
+        return Err(BenchError::Config(
+            "Bucket config must have at least one entry".into(),
+        ));
+    }
+
+    Ok(buckets)
+}
+
+/// JSON fragment prefix/suffix for image content blocks.
+const IMG_JSON_PREFIX: &str = r#"{"type":"image_url","image_url":{"url":"data:image/jpeg;base64,"#;
+const IMG_JSON_SUFFIX: &str = r#""}}"#;
+
+/// Generate a synthetic random JPEG image and return it as a pre-serialized JSON fragment.
+///
+/// Builds the complete JSON string in a single allocation:
+/// `{"type":"image_url","image_url":{"url":"data:image/jpeg;base64,<b64>"}}`
+///
+/// The base64 data is written directly into the final string — no intermediate
+/// String or format!() copy.
+fn generate_random_image(width: u32, height: u32, rng: &mut StdRng) -> Result<Arc<str>> {
+    let pixel_count = (width as usize) * (height as usize) * 3;
+    let mut pixels = vec![0u8; pixel_count];
+    rng.fill(pixels.as_mut_slice());
+
+    let img = image::RgbImage::from_raw(width, height, pixels)
+        .ok_or_else(|| BenchError::Config("Failed to create image from random pixels".into()))?;
+
+    // Pre-allocate JPEG buffer (random pixels compress poorly, estimate ~60% of raw)
+    let estimated_jpeg = pixel_count * 3 / 5;
+    let mut buf = Cursor::new(Vec::with_capacity(estimated_jpeg));
+    img.write_to(&mut buf, image::ImageFormat::Jpeg)
+        .map_err(|e| BenchError::Config(format!("Failed to encode JPEG: {e}")))?;
+
+    let jpeg_bytes = buf.into_inner();
+
+    // Pre-compute exact output size: prefix + base64_len + suffix
+    let b64_len = jpeg_bytes.len().div_ceil(3) * 4;
+    let total_len = IMG_JSON_PREFIX.len() + b64_len + IMG_JSON_SUFFIX.len();
+
+    // Single allocation: write base64 directly into the JSON fragment string
+    let mut json_fragment = String::with_capacity(total_len);
+    json_fragment.push_str(IMG_JSON_PREFIX);
+    base64::engine::general_purpose::STANDARD.encode_string(&jpeg_bytes, &mut json_fragment);
+    json_fragment.push_str(IMG_JSON_SUFFIX);
+
+    Ok(Arc::from(json_fragment))
+}
+
+/// Sample multimodal items for a single request.
+///
+/// Returns a list of (height, width, num_frames) tuples.
+fn sample_mm_items(
+    rng: &mut StdRng,
+    min_items: usize,
+    max_items: usize,
+    buckets: &[(MmBucketKey, f64)],
+    limit: &MmLimitPerPrompt,
+) -> Vec<MmBucketKey> {
+    let num_items = if min_items == max_items {
+        min_items
+    } else {
+        rng.random_range(min_items..=max_items)
+    };
+
+    // Filter to non-zero probability buckets
+    let active_buckets: Vec<&(MmBucketKey, f64)> =
+        buckets.iter().filter(|(_, p)| *p > 0.0).collect();
+    if active_buckets.is_empty() || num_items == 0 {
+        return Vec::new();
+    }
+
+    let total_weight: f64 = active_buckets.iter().map(|(_, p)| p).sum();
+    if total_weight <= 0.0 {
+        return Vec::new();
+    }
+
+    let mut result = Vec::with_capacity(num_items);
+    let mut image_count = 0usize;
+    let mut video_count = 0usize;
+
+    for _ in 0..num_items {
+        // Build normalized weights considering remaining capacity
+        let mut weights: Vec<f64> = Vec::with_capacity(active_buckets.len());
+        for (key, prob) in &active_buckets {
+            let is_video = key.num_frames > 1;
+            let at_limit = if is_video {
+                video_count >= limit.video
+            } else {
+                image_count >= limit.image
+            };
+            weights.push(if at_limit { 0.0 } else { *prob });
+        }
+
+        let w_total: f64 = weights.iter().sum();
+        if w_total <= 0.0 {
+            break; // All modalities at limit
+        }
+
+        // Weighted random selection (strict `<` to avoid selecting zero-weight buckets)
+        let r = rng.random::<f64>() * w_total;
+        let mut cumulative = 0.0;
+        // Default to last non-zero-weight bucket (floating-point accumulation fallback)
+        let mut selected_idx = weights.iter().rposition(|w| *w > 0.0).unwrap_or(0);
+        for (i, w) in weights.iter().enumerate() {
+            cumulative += w;
+            if r < cumulative {
+                selected_idx = i;
+                break;
+            }
+        }
+
+        let (key, _) = &active_buckets[selected_idx];
+        if key.num_frames > 1 {
+            video_count += 1;
+        } else {
+            image_count += 1;
+        }
+        result.push(key.clone());
+    }
+
+    result
+}
+
+/// Generate random multimodal dataset.
+///
+/// Mirrors Python's RandomMultiModalDataset.sample() from datasets.py.
+/// Generates text prompts with exact token lengths and random images/videos.
+pub fn generate_random_mm_dataset(
+    tokenizer: &TokenizerKind,
+    num_requests: usize,
+    input_len: usize,
+    output_len: usize,
+    prefix_len: usize,
+    range_ratio: crate::config::RangeRatio,
+    seed: u64,
+    request_id_prefix: &str,
+    base_items_per_request: usize,
+    num_mm_items_range_ratio: f64,
+    limit: &MmLimitPerPrompt,
+    buckets: &[(MmBucketKey, f64)],
+    enable_multimodal_chat: bool,
+) -> Result<Vec<SampleRequest>> {
+    if !(0.0..=1.0).contains(&num_mm_items_range_ratio) {
+        return Err(BenchError::Config(
+            "num_mm_items_range_ratio must be in [0, 1]".into(),
+        ));
+    }
+
+    // Check for video buckets with non-zero probability
+    for (key, prob) in buckets {
+        if key.num_frames > 1 && *prob > 0.0 {
+            return Err(BenchError::Config(
+                "Video generation (num_frames > 1) is not yet supported in Rust. \
+                 Set video bucket probabilities to 0.0."
+                    .into(),
+            ));
+        }
+    }
+
+    // Compute item count bounds
+    let n = base_items_per_request as f64;
+    let r = num_mm_items_range_ratio;
+    let min_items = (n * (1.0 - r)).floor().max(0.0) as usize;
+    let max_items = (n * (1.0 + r)).ceil() as usize;
+    // Clamp to total modality limit
+    let total_limit = limit.image + limit.video;
+    let max_items = max_items.min(total_limit);
+    let min_items = min_items.min(max_items);
+
+    let vocab_size = tokenizer.vocab_size();
+    let allowed_tokens = tokenizer.get_allowed_tokens();
+    if allowed_tokens.is_empty() {
+        return Err(BenchError::Tokenizer("No allowed tokens found".into()));
+    }
+
+    let num_special = tokenizer.num_special_tokens_to_add();
+    let real_input_len = input_len.saturating_sub(num_special);
+
+    // Python semantics: sample uniformly from [len*(1-r), len*(1+r)].
+    let (input_low, input_high) = range_ratio.input_bounds(real_input_len);
+    let (output_low, output_high) = range_ratio.output_bounds(output_len);
+
+    // Pre-generate per-request params
+    let mut rng = StdRng::seed_from_u64(seed);
+
+    struct RequestParams {
+        input_len: usize,
+        output_len: usize,
+        offset: usize,
+    }
+
+    let params: Vec<RequestParams> = (0..num_requests)
+        .map(|_| {
+            let il = if input_low == input_high {
+                input_low
+            } else {
+                rng.random_range(input_low..=input_high)
+            };
+            let ol = if output_low == output_high {
+                output_low
+            } else {
+                rng.random_range(output_low..=output_high)
+            };
+            let off = rng.random_range(0..vocab_size as usize);
+            RequestParams {
+                input_len: il,
+                output_len: ol,
+                offset: off,
+            }
+        })
+        .collect();
+
+    // Pre-generate multimodal item configs per request
+    let mm_configs: Vec<Vec<MmBucketKey>> = (0..num_requests)
+        .map(|_| sample_mm_items(&mut rng, min_items, max_items, buckets, limit))
+        .collect();
+
+    // Generate text prompts (need text for chat backend, not just token IDs)
+    let prefix_token_ids = if prefix_len > 0 {
+        generate_prefix(tokenizer, &allowed_tokens, prefix_len, seed)?
+    } else {
+        Vec::new()
+    };
+
+    // Generate token sequences
+    let prefix_ref = &prefix_token_ids;
+    let allowed_ref = &allowed_tokens;
+
+    let token_sequences: Vec<Vec<u32>> = params
+        .par_iter()
+        .enumerate()
+        .map(|(i, p)| {
+            let at_len = allowed_ref.len();
+            let mut seq = Vec::with_capacity(prefix_ref.len() + p.input_len);
+            seq.extend_from_slice(prefix_ref);
+            for j in 0..p.input_len {
+                seq.push(allowed_ref[(p.offset + i + j) % at_len]);
+            }
+            seq
+        })
+        .collect();
+
+    let target_lens: Vec<usize> = params.iter().map(|p| prefix_len + p.input_len).collect();
+
+    // Decode tokens to text (chat backend needs text prompts for multimodal)
+    let prompts: Result<Vec<String>> = token_sequences
+        .into_par_iter()
+        .enumerate()
+        .map(|(i, tokens)| {
+            let (text, _adjusted) = super::random::gen_prompt_decode_to_target_len(
+                tokenizer,
+                &tokens,
+                target_lens[i],
+                false,
+                allowed_ref,
+            )?;
+            Ok(text)
+        })
+        .collect();
+    let prompts = prompts?;
+
+    // Generate images for each request (parallel per request)
+    // Each request gets its own RNG seeded deterministically.
+    let rid_prefix = request_id_prefix.to_string();
+    let result: Vec<SampleRequest> = prompts
+        .into_par_iter()
+        .enumerate()
+        .map(|(i, prompt)| {
+            let mut item_rng =
+                StdRng::seed_from_u64(seed.wrapping_add(i as u64).wrapping_add(0xBEEF));
+            let mm_items: Vec<Arc<str>> = mm_configs[i]
+                .iter()
+                .map(|key| {
+                    generate_random_image(key.width, key.height, &mut item_rng)
+                        .expect("Image generation should not fail")
+                })
+                .collect();
+
+            let mm_content: Option<Arc<[Arc<str>]>> = if mm_items.is_empty() {
+                None
+            } else {
+                Some(Arc::from(mm_items))
+            };
+
+            // --enable-multimodal-chat: pre-build the full chat `messages` array
+            // (text part + mm items) at dataset time, mirroring Python's
+            // apply_multimodal_chat_transformation. mm content moves inside the
+            // messages string; the backend splices it verbatim.
+            let (mm_content, chat_messages_json) = if enable_multimodal_chat {
+                let msgs = build_chat_messages_json(&prompt, mm_content.as_deref());
+                (None, Some(Arc::from(msgs.as_str())))
+            } else {
+                (mm_content, None)
+            };
+
+            SampleRequest {
+                prompt: Arc::from(prompt.as_str()),
+                prompt_len: target_lens[i],
+                expected_output_len: params[i].output_len,
+                request_id: Some(format!("{rid_prefix}{i}")),
+                multi_modal_content: mm_content,
+                chat_messages_json,
+                ..Default::default()
+            }
+        })
+        .collect();
+
+    Ok(result)
+}
+
+/// Pre-serialize the OpenAI chat `messages` array for --enable-multimodal-chat.
+///
+/// Produces `[{"role":"user","content":[{"type":"text","text":"..."},<frag>,...]}]`
+/// by concatenating the JSON-escaped prompt with the pre-serialized mm fragments,
+/// so the ~200KB+ base64 image data is never parsed or re-serialized.
+pub(crate) fn build_chat_messages_json(prompt: &str, mm_items: Option<&[Arc<str>]>) -> String {
+    let mm_total: usize =
+        mm_items.map(|items| items.iter().map(|f| f.len() + 1).sum()).unwrap_or(0);
+    let mut msgs = String::with_capacity(64 + prompt.len() * 2 + mm_total);
+    msgs.push_str(r#"[{"role":"user","content":[{"type":"text","text":"#);
+    // serde_json::to_string on &str produces a JSON-escaped quoted string
+    msgs.push_str(&serde_json::to_string(prompt).unwrap());
+    msgs.push('}');
+    for fragment in mm_items.unwrap_or(&[]) {
+        msgs.push(',');
+        msgs.push_str(fragment);
+    }
+    msgs.push_str("]}]");
+    msgs
+}
+
+fn generate_prefix(
+    tokenizer: &TokenizerKind,
+    allowed_tokens: &[u32],
+    prefix_len: usize,
+    seed: u64,
+) -> Result<Vec<u32>> {
+    let mut rng = StdRng::seed_from_u64(seed.wrapping_add(0xDEAD));
+    let tokens: Vec<u32> = (0..prefix_len)
+        .map(|_| allowed_tokens[rng.random_range(0..allowed_tokens.len())])
+        .collect();
+
+    let (_, adjusted) = super::random::gen_prompt_decode_to_target_len(
+        tokenizer,
+        &tokens,
+        prefix_len,
+        false,
+        allowed_tokens,
+    )?;
+    Ok(adjusted)
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+
+    #[test]
+    fn test_parse_bucket_config_basic() {
+        let input = "{(256,256,1): 0.5, (720,1280,1): 0.5}";
+        let buckets = parse_bucket_config(input).unwrap();
+        assert_eq!(buckets.len(), 2);
+        assert_eq!(buckets[0].0.height, 256);
+        assert_eq!(buckets[0].0.width, 256);
+        assert_eq!(buckets[0].0.num_frames, 1);
+        assert!((buckets[0].1 - 0.5).abs() < 1e-10);
+        assert_eq!(buckets[1].0.height, 720);
+        assert_eq!(buckets[1].0.width, 1280);
+    }
+
+    #[test]
+    fn test_parse_bucket_config_single() {
+        let input = "{(1024, 800, 1): 1.0}";
+        let buckets = parse_bucket_config(input).unwrap();
+        assert_eq!(buckets.len(), 1);
+        assert_eq!(buckets[0].0.height, 1024);
+        assert_eq!(buckets[0].0.width, 800);
+        assert_eq!(buckets[0].0.num_frames, 1);
+        assert!((buckets[0].1 - 1.0).abs() < 1e-10);
+    }
+
+    #[test]
+    fn test_parse_bucket_config_with_video() {
+        let input = "{(256,256,1): 0.4, (720,1280,1): 0.4, (720,1280,16): 0.2}";
+        let buckets = parse_bucket_config(input).unwrap();
+        assert_eq!(buckets.len(), 3);
+        assert_eq!(buckets[2].0.num_frames, 16);
+    }
+
+    #[test]
+    fn test_build_chat_messages_json_valid_and_ordered() {
+        let frag: Arc<str> =
+            Arc::from(r#"{"type":"image_url","image_url":{"url":"data:image/jpeg;base64,AAAA"}}"#);
+        let msgs = build_chat_messages_json("hi \"there\"\nline2", Some(&[frag]));
+        let v: serde_json::Value = serde_json::from_str(&msgs).expect("must be valid JSON");
+        assert_eq!(v.as_array().unwrap().len(), 1);
+        assert_eq!(v[0]["role"], "user");
+        let content = v[0]["content"].as_array().unwrap();
+        assert_eq!(content.len(), 2);
+        assert_eq!(content[0]["type"], "text");
+        assert_eq!(content[0]["text"], "hi \"there\"\nline2");
+        assert_eq!(content[1]["type"], "image_url");
+    }
+
+    #[test]
+    fn test_build_chat_messages_json_text_only() {
+        let msgs = build_chat_messages_json("plain", None);
+        let v: serde_json::Value = serde_json::from_str(&msgs).unwrap();
+        let content = v[0]["content"].as_array().unwrap();
+        assert_eq!(content.len(), 1);
+        assert_eq!(content[0]["text"], "plain");
+    }
+
+    #[test]
+    fn test_parse_limit_mm_per_prompt() {
+        let input = r#"{"image": 3, "video": 0}"#;
+        let limit = parse_limit_mm_per_prompt(input).unwrap();
+        assert_eq!(limit.image, 3);
+        assert_eq!(limit.video, 0);
+    }
+
+    #[test]
+    fn test_parse_limit_mm_per_prompt_defaults() {
+        let input = r#"{}"#;
+        let limit = parse_limit_mm_per_prompt(input).unwrap();
+        assert_eq!(limit.image, 255);
+        assert_eq!(limit.video, 1);
+    }
+
+    #[test]
+    fn test_sample_mm_items_basic() {
+        let mut rng = StdRng::seed_from_u64(42);
+        let buckets = vec![
+            (
+                MmBucketKey {
+                    height: 256,
+                    width: 256,
+                    num_frames: 1,
+                },
+                0.5,
+            ),
+            (
+                MmBucketKey {
+                    height: 720,
+                    width: 1280,
+                    num_frames: 1,
+                },
+                0.5,
+            ),
+        ];
+        let limit = MmLimitPerPrompt { image: 5, video: 0 };
+        let items = sample_mm_items(&mut rng, 2, 3, &buckets, &limit);
+        assert!(items.len() >= 2 && items.len() <= 3);
+        for item in &items {
+            assert_eq!(item.num_frames, 1);
+        }
+    }
+
+    #[test]
+    fn test_sample_mm_items_respects_limit() {
+        let mut rng = StdRng::seed_from_u64(42);
+        let buckets = vec![(
+            MmBucketKey {
+                height: 256,
+                width: 256,
+                num_frames: 1,
+            },
+            1.0,
+        )];
+        let limit = MmLimitPerPrompt { image: 2, video: 0 };
+        let items = sample_mm_items(&mut rng, 5, 5, &buckets, &limit);
+        // Should be capped at 2 due to image limit
+        assert_eq!(items.len(), 2);
+    }
+
+    #[test]
+    fn test_generate_random_image() {
+        let mut rng = StdRng::seed_from_u64(42);
+        let result = generate_random_image(64, 64, &mut rng).unwrap();
+        // Result is a pre-serialized JSON fragment
+        assert!(
+            result
+                .starts_with(r#"{"type":"image_url","image_url":{"url":"data:image/jpeg;base64,"#)
+        );
+        assert!(result.ends_with(r#""}}"#));
+        // Verify it's valid JSON
+        let parsed: serde_json::Value = serde_json::from_str(&result).unwrap();
+        assert_eq!(parsed["type"], "image_url");
+    }
+}
@@ -0,0 +1,185 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+//! Random dataset specialized for scoring/rerank benchmarks: each request is
+//! one query plus a batch of documents. Mirrors Python's
+//! `RandomDatasetForReranking`.
+//!
+//! With `is_reranker` (default): the query and each document share the
+//! request's token budget (`query + sep + doc ~= input_len`), and every
+//! batched request counts the query once per document pair.
+//! With `--no-reranker` (embedding-based scoring): the query is just another
+//! embedding input occupying the first batch slot.
+
+use std::sync::Arc;
+
+use rand::rngs::StdRng;
+use rand::{Rng, SeedableRng};
+use rayon::prelude::*;
+
+use super::SampleRequest;
+use crate::config::RangeRatio;
+use crate::error::{BenchError, Result};
+use crate::tokenizer::TokenizerKind;
+
+pub fn generate_random_rerank_dataset(
+    tokenizer: &TokenizerKind,
+    num_requests: usize,
+    input_len: usize,
+    range_ratio: RangeRatio,
+    seed: u64,
+    request_id_prefix: &str,
+    batch_size: usize,
+    is_reranker: bool,
+) -> Result<Vec<SampleRequest>> {
+    let allowed_tokens = tokenizer.get_allowed_tokens();
+    if allowed_tokens.is_empty() {
+        return Err(BenchError::Tokenizer("No allowed tokens found".into()));
+    }
+    let allowed_ref = &allowed_tokens;
+
+    let num_special = tokenizer.num_special_tokens_to_add();
+    let real_input_len = input_len.saturating_sub(num_special);
+
+    let n_sep_tokens = usize::from(is_reranker);
+    let query_len_param = if is_reranker {
+        (real_input_len / 2).saturating_sub(n_sep_tokens)
+    } else {
+        real_input_len
+    };
+
+    let mut rng = StdRng::seed_from_u64(seed);
+    let sample = |rng: &mut StdRng, (low, high): (usize, usize)| -> usize {
+        if low == high {
+            low
+        } else {
+            rng.random_range(low..=high)
+        }
+    };
+
+    // One query length for the whole run, like Python.
+    let query_len = sample(&mut rng, range_ratio.input_bounds(query_len_param));
+
+    // --no-reranker folds the query into the first batch slot.
+    let (num_docs, docs_per_batch, doc_len_param) = if is_reranker {
+        let doc_len = real_input_len.saturating_sub(query_len).saturating_sub(n_sep_tokens);
+        (num_requests, batch_size, doc_len)
+    } else {
+        (num_requests - 1, batch_size - 1, real_input_len)
+    };
+    if doc_len_param == 0 {
+        return Err(BenchError::Config(format!(
+            "random-rerank: --random-input-len {input_len} leaves no budget for documents \
+             (query_len={query_len})"
+        )));
+    }
+
+    // Pre-sample per-document lengths and offsets deterministically.
+    let doc_bounds = range_ratio.input_bounds(doc_len_param);
+    let doc_params: Vec<(usize, usize)> = (0..num_docs)
+        .map(|_| {
+            (
+                sample(&mut rng, doc_bounds),
+                rng.random_range(0..allowed_ref.len()),
+            )
+        })
+        .collect();
+    let query_offset = rng.random_range(0..allowed_ref.len());
+
+    // Exact-length text: token sequence -> decode -> re-encode -> truncate -> decode.
+    let gen_text = |target: usize, offset: usize, index: usize| -> Result<(Arc<str>, usize)> {
+        let at_len = allowed_ref.len();
+        let tokens: Vec<u32> =
+            (0..target).map(|j| allowed_ref[(offset + index + j) % at_len]).collect();
+        let text = tokenizer.decode(&tokens, true)?;
+        let mut re_encoded = tokenizer.encode(&text, false)?;
+        re_encoded.truncate(target);
+        let final_text = tokenizer.decode(&re_encoded, true)?;
+        Ok((Arc::from(final_text), re_encoded.len()))
+    };
+
+    let (query_prompt, query_input_len) = gen_text(query_len, query_offset, 0)?;
+
+    let docs: Vec<(Arc<str>, usize)> = doc_params
+        .par_iter()
+        .enumerate()
+        .map(|(i, (len, offset))| gen_text(*len, *offset, i + 1))
+        .collect::<Result<Vec<_>>>()?;
+
+    // Batch documents; every request is [query, doc1, doc2, ...].
+    let rid_prefix = request_id_prefix.to_string();
+    let requests = docs
+        .chunks(docs_per_batch)
+        .enumerate()
+        .map(|(batch_idx, batch)| {
+            let query_contrib = if is_reranker {
+                (query_input_len + n_sep_tokens) * batch.len()
+            } else {
+                query_input_len
+            };
+            let mut prompt_list: Vec<Arc<str>> = Vec::with_capacity(batch.len() + 1);
+            prompt_list.push(query_prompt.clone());
+            prompt_list.extend(batch.iter().map(|(text, _)| text.clone()));
+            SampleRequest {
+                prompt_list: Some(Arc::from(prompt_list)),
+                prompt_len: query_contrib + batch.iter().map(|(_, len)| len).sum::<usize>(),
+                expected_output_len: 0,
+                request_id: Some(format!("{rid_prefix}{batch_idx}")),
+                ..Default::default()
+            }
+        })
+        .collect();
+
+    Ok(requests)
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+
+    /// gpt2 via built-in tiktoken encoding — loads without network access.
+    fn test_tokenizer() -> TokenizerKind {
+        crate::tokenizer::load_tokenizer("gpt2", false, None)
+            .expect("gpt2 built-in tiktoken should always load without network")
+    }
+
+    fn fixed_ratio() -> RangeRatio {
+        RangeRatio::parse("0.0").unwrap()
+    }
+
+    #[test]
+    fn test_random_rerank_reranker_mode() {
+        let tok = test_tokenizer();
+        // 6 docs in batches of 3 -> 2 requests of [query, d1, d2, d3]
+        let reqs = generate_random_rerank_dataset(&tok, 6, 128, fixed_ratio(), 0, "t-", 3, true)
+            .expect("generation should succeed");
+        assert_eq!(reqs.len(), 2);
+        for r in &reqs {
+            let list = r.prompt_list.as_ref().expect("prompt_list must be set");
+            assert_eq!(list.len(), 4);
+            assert_eq!(r.expected_output_len, 0);
+            assert!(r.prompt_len > 0);
+        }
+        // Same query shared across requests
+        assert_eq!(
+            reqs[0].prompt_list.as_ref().unwrap()[0],
+            reqs[1].prompt_list.as_ref().unwrap()[0]
+        );
+        // Reranker budget: query+sep+doc pairs stay near input_len per pair
+        // (query ~63, doc ~64 for input_len=128, gpt2 has no special tokens)
+        let list = reqs[0].prompt_list.as_ref().unwrap();
+        let query_tokens = tok.encode(&list[0], false).unwrap().len();
+        assert!(query_tokens <= 64, "query too long: {query_tokens}");
+    }
+
+    #[test]
+    fn test_random_rerank_no_reranker_mode() {
+        let tok = test_tokenizer();
+        // no-reranker: query occupies first slot; 5 non-query docs in batches of 2
+        let reqs = generate_random_rerank_dataset(&tok, 6, 64, fixed_ratio(), 0, "t-", 3, false)
+            .expect("generation should succeed");
+        // 6-1=5 docs, batches of 3-1=2 -> 3 requests
+        assert_eq!(reqs.len(), 3);
+        assert_eq!(reqs[0].prompt_list.as_ref().unwrap().len(), 3);
+    }
+}
@@ -0,0 +1,209 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+use std::sync::Arc;
+
+use rand::rngs::StdRng;
+use rand::seq::SliceRandom;
+use rand::{Rng, SeedableRng};
+
+use super::SampleRequest;
+use crate::error::{BenchError, Result};
+use crate::tokenizer::TokenizerKind;
+
+/// Default validation bounds matching Python's is_valid_sequence() defaults.
+const MIN_LEN: usize = 4;
+const MAX_PROMPT_LEN: usize = 1024;
+const MAX_TOTAL_LEN: usize = 2048;
+
+/// Default HuggingFace dataset repo and filename for ShareGPT.
+const DEFAULT_SHAREGPT_REPO: &str = "anon8231489123/ShareGPT_Vicuna_unfiltered";
+const DEFAULT_SHAREGPT_FILE: &str = "ShareGPT_V3_unfiltered_cleaned_split.json";
+
+/// Download the default ShareGPT dataset from HuggingFace Hub.
+/// Uses hf-hub's built-in cache — subsequent calls return the cached path instantly.
+pub fn download_sharegpt_dataset() -> Result<String> {
+    println!(
+        "Downloading ShareGPT dataset from {DEFAULT_SHAREGPT_REPO}/{DEFAULT_SHAREGPT_FILE} ..."
+    );
+    let repo = crate::hub::HubRepo::dataset(DEFAULT_SHAREGPT_REPO.to_string());
+    let path = repo.get(DEFAULT_SHAREGPT_FILE).map_err(|e| {
+        BenchError::Config(format!(
+            "Failed to download ShareGPT dataset from '{DEFAULT_SHAREGPT_REPO}': {e}"
+        ))
+    })?;
+    let path_str = path.to_string_lossy().to_string();
+    println!("ShareGPT dataset ready: {path_str}");
+    Ok(path_str)
+}
+
+/// Load and sample from a ShareGPT-format JSON dataset.
+///
+/// Mirrors Python's ShareGPTDataset from datasets.py:1230-1313.
+pub fn load_sharegpt_dataset(
+    tokenizer: &TokenizerKind,
+    dataset_path: &str,
+    num_requests: usize,
+    output_len_override: Option<usize>,
+    seed: u64,
+    request_id_prefix: &str,
+    no_oversample: bool,
+    disable_shuffle: bool,
+) -> Result<Vec<SampleRequest>> {
+    // Load JSON file
+    let content = std::fs::read_to_string(dataset_path).map_err(|e| {
+        BenchError::Config(format!(
+            "Failed to read ShareGPT file '{dataset_path}': {e}"
+        ))
+    })?;
+
+    let data: serde_json::Value = serde_json::from_str(&content)
+        .map_err(|e| BenchError::Config(format!("Invalid JSON in ShareGPT file: {e}")))?;
+
+    let entries = data
+        .as_array()
+        .ok_or_else(|| BenchError::Config("ShareGPT file must contain a JSON array".into()))?;
+
+    // Filter entries with at least 2 conversation turns
+    let mut filtered: Vec<&serde_json::Value> = entries
+        .iter()
+        .filter(|entry| {
+            entry
+                .get("conversations")
+                .and_then(|c| c.as_array())
+                .map(|a| a.len() >= 2)
+                .unwrap_or(false)
+        })
+        .collect();
+
+    if filtered.is_empty() {
+        return Err(BenchError::Config(
+            "No valid entries in ShareGPT file (need at least 2 conversation turns)".into(),
+        ));
+    }
+
+    // Shuffle (unless disabled)
+    let mut rng = StdRng::seed_from_u64(seed);
+    if !disable_shuffle {
+        filtered.shuffle(&mut rng);
+    }
+
+    // Sample requests
+    let mut samples = Vec::new();
+    let mut ind = 0;
+
+    for entry in &filtered {
+        if samples.len() >= num_requests {
+            break;
+        }
+
+        let conversations = entry["conversations"].as_array().unwrap();
+        let prompt = conversations[0]["value"].as_str().unwrap_or("");
+        let completion = conversations[1]["value"].as_str().unwrap_or("");
+
+        if prompt.is_empty() {
+            continue;
+        }
+
+        // Tokenize prompt and completion
+        let prompt_ids = tokenizer.encode(prompt, false)?;
+        let prompt_len = prompt_ids.len();
+
+        let new_output_len = if let Some(override_len) = output_len_override {
+            override_len
+        } else {
+            let completion_ids = tokenizer.encode(completion, false)?;
+            completion_ids.len()
+        };
+
+        // Validate sequence lengths (matching Python's is_valid_sequence)
+        let skip_min_output = output_len_override.is_some();
+        if !is_valid_sequence(prompt_len, new_output_len, skip_min_output) {
+            continue;
+        }
+
+        samples.push(SampleRequest {
+            prompt: Arc::from(prompt),
+            prompt_len,
+            expected_output_len: new_output_len,
+            request_id: Some(format!("{request_id_prefix}{ind}")),
+            ..Default::default()
+        });
+        ind += 1;
+    }
+
+    // Oversample if dataset is smaller than requested
+    if samples.len() < num_requests {
+        if no_oversample {
+            println!(
+                "Skipping oversampling. Total samples: {} (requested: {num_requests})",
+                samples.len()
+            );
+        } else if !samples.is_empty() {
+            let needed = num_requests - samples.len();
+            let original_len = samples.len();
+            for i in 0..needed {
+                let mut req = samples[rng.random_range(0..original_len)].clone();
+                req.request_id = Some(format!("{request_id_prefix}{}", original_len + i));
+                samples.push(req);
+            }
+            println!(
+                "Oversampled requests from {original_len} to {} total samples.",
+                samples.len()
+            );
+        }
+    }
+
+    if samples.is_empty() {
+        return Err(BenchError::Config(
+            "No valid samples after filtering ShareGPT dataset. \
+             Try relaxing constraints or using a larger dataset."
+                .into(),
+        ));
+    }
+
+    Ok(samples)
+}
+
+/// Validate a sequence based on prompt and output lengths.
+/// Mirrors Python's is_valid_sequence() from datasets.py:260-284.
+fn is_valid_sequence(
+    prompt_len: usize,
+    output_len: usize,
+    skip_min_output_len_check: bool,
+) -> bool {
+    if prompt_len < MIN_LEN {
+        return false;
+    }
+    if !skip_min_output_len_check && output_len < MIN_LEN {
+        return false;
+    }
+    if prompt_len > MAX_PROMPT_LEN {
+        return false;
+    }
+    if prompt_len + output_len > MAX_TOTAL_LEN {
+        return false;
+    }
+    true
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+
+    #[test]
+    fn test_is_valid_sequence() {
+        // Valid
+        assert!(is_valid_sequence(100, 50, false));
+        // Prompt too short
+        assert!(!is_valid_sequence(3, 50, false));
+        // Output too short
+        assert!(!is_valid_sequence(100, 3, false));
+        // Output too short but skip check
+        assert!(is_valid_sequence(100, 1, true));
+        // Prompt too long
+        assert!(!is_valid_sequence(1025, 50, false));
+        // Combined too long
+        assert!(!is_valid_sequence(1024, 1025, false));
+    }
+}
@@ -0,0 +1,131 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+use std::sync::Arc;
+
+use rand::SeedableRng;
+use rand::rngs::StdRng;
+use rand::seq::IndexedRandom;
+
+use super::SampleRequest;
+use crate::error::{BenchError, Result};
+use crate::tokenizer::TokenizerKind;
+
+/// Default values mirror Python's SonnetDataset defaults (datasets.py).
+pub const DEFAULT_PREFIX_LEN: usize = 200;
+pub const DEFAULT_INPUT_LEN: usize = 550;
+pub const DEFAULT_OUTPUT_LEN: usize = 150;
+
+const BASE_PROMPT: &str = "Pick as many lines as you can from these poem lines:\n";
+
+/// Shakespeare's sonnets, public domain. Bundled so `--dataset-name sonnet` works
+/// out of the box without `--dataset-path`. Source:
+/// https://raw.githubusercontent.com/vllm-project/vllm/main/benchmarks/sonnet.txt
+const BUILTIN_SONNET: &str = include_str!("sonnet.txt");
+
+/// Load the sonnet dataset and generate `num_requests` prompts targeting `input_len`
+/// total prompt tokens.
+///
+/// Mirrors Python's `SonnetDataset.sample()` from `vllm/benchmarks/datasets/datasets.py`.
+/// The Rust port skips `apply_chat_template` (no Jinja runtime here): `base_offset`
+/// and `prompt_len` are computed from the raw text. The resulting prompts are slightly
+/// shorter than Python's chat-template-formatted version (off by the chat scaffolding
+/// tokens, typically <20).
+pub fn load_sonnet_dataset(
+    tokenizer: &TokenizerKind,
+    dataset_path: Option<&str>,
+    num_requests: usize,
+    input_len: usize,
+    output_len: usize,
+    prefix_len: usize,
+    seed: u64,
+    request_id_prefix: &str,
+) -> Result<Vec<SampleRequest>> {
+    let content = match dataset_path {
+        Some(path) => std::fs::read_to_string(path)
+            .map_err(|e| BenchError::Config(format!("Failed to read sonnet file '{path}': {e}")))?,
+        None => BUILTIN_SONNET.to_string(),
+    };
+
+    // Match Python's f.readlines(): keep trailing newlines so that joining lines
+    // reconstructs the original text without inserting extra separators.
+    let lines: Vec<String> = content.split_inclusive('\n').map(|s| s.to_string()).collect();
+
+    if lines.is_empty() {
+        let src = dataset_path.unwrap_or("<built-in>");
+        return Err(BenchError::Config(format!("Sonnet file '{src}' is empty")));
+    }
+
+    // Average tokens per line (used to estimate how many lines to draw).
+    let mut total_tokens: usize = 0;
+    for line in &lines {
+        let ids = tokenizer.encode(line, false)?;
+        total_tokens += ids.len();
+    }
+    let avg_len = total_tokens as f64 / lines.len() as f64;
+    if avg_len <= 0.0 {
+        return Err(BenchError::Config(
+            "Sonnet lines tokenized to zero tokens on average".into(),
+        ));
+    }
+
+    let base_ids = tokenizer.encode(BASE_PROMPT, false)?;
+    let base_offset = base_ids.len();
+    if input_len <= base_offset {
+        return Err(BenchError::Config(format!(
+            "--sonnet-input-len ({input_len}) must be larger than the base prompt length ({base_offset})"
+        )));
+    }
+
+    let num_input_lines = ((input_len - base_offset) as f64 / avg_len).round() as i64;
+    let num_prefix_lines =
+        (((prefix_len as i64 - base_offset as i64) as f64) / avg_len).round().max(0.0) as i64;
+    let num_input_lines = num_input_lines.max(0) as usize;
+    let num_prefix_lines = (num_prefix_lines as usize).min(lines.len());
+    let num_input_lines = num_input_lines.max(num_prefix_lines);
+
+    let prefix_lines: &[String] = &lines[..num_prefix_lines];
+    let extras_per_request = num_input_lines - num_prefix_lines;
+
+    let mut rng = StdRng::seed_from_u64(seed);
+    let mut samples = Vec::with_capacity(num_requests);
+    let mut ind = 0usize;
+    let mut attempts = 0usize;
+    let max_attempts = num_requests.saturating_mul(20).max(1000);
+
+    while samples.len() < num_requests {
+        if attempts >= max_attempts {
+            return Err(BenchError::Config(format!(
+                "Could not assemble {num_requests} sonnet prompts under input_len={input_len} \
+                 after {attempts} attempts. Try increasing --sonnet-input-len."
+            )));
+        }
+        attempts += 1;
+
+        let mut prompt = String::with_capacity(BASE_PROMPT.len() + 256 * num_input_lines);
+        prompt.push_str(BASE_PROMPT);
+        for line in prefix_lines {
+            prompt.push_str(line);
+        }
+        for _ in 0..extras_per_request {
+            // random.choices with replacement — duplicates are allowed.
+            let line = lines.choose(&mut rng).unwrap();
+            prompt.push_str(line);
+        }
+
+        let prompt_ids = tokenizer.encode(&prompt, false)?;
+        let prompt_len = prompt_ids.len();
+        if prompt_len <= input_len {
+            samples.push(SampleRequest {
+                prompt: Arc::from(prompt),
+                prompt_len,
+                expected_output_len: output_len,
+                request_id: Some(format!("{request_id_prefix}{ind}")),
+                ..Default::default()
+            });
+            ind += 1;
+        }
+    }
+
+    Ok(samples)
+}
@@ -0,0 +1,518 @@
+FROM fairest creatures we desire increase,
+That thereby beauty's rose might never die,
+But as the riper should by time decease,
+His tender heir might bear his memory:
+But thou, contracted to thine own bright eyes,
+Feed'st thy light'st flame with self-substantial fuel,
+Making a famine where abundance lies,
+Thyself thy foe, to thy sweet self too cruel.
+Thou that art now the world's fresh ornament
+And only herald to the gaudy spring,
+Within thine own bud buriest thy content
+And, tender churl, makest waste in niggarding.
+Pity the world, or else this glutton be,
+To eat the world's due, by the grave and thee.
+When forty winters shall beseige thy brow,
+And dig deep trenches in thy beauty's field,
+Thy youth's proud livery, so gazed on now,
+Will be a tatter'd weed, of small worth held:
+Then being ask'd where all thy beauty lies,
+Where all the treasure of thy lusty days,
+To say, within thine own deep-sunken eyes,
+Were an all-eating shame and thriftless praise.
+How much more praise deserved thy beauty's use,
+If thou couldst answer 'This fair child of mine
+Shall sum my count and make my old excuse,'
+Proving his beauty by succession thine!
+This were to be new made when thou art old,
+And see thy blood warm when thou feel'st it cold.
+Look in thy glass, and tell the face thou viewest
+Now is the time that face should form another;
+Whose fresh repair if now thou not renewest,
+Thou dost beguile the world, unbless some mother.
+For where is she so fair whose unear'd womb
+Disdains the tillage of thy husbandry?
+Or who is he so fond will be the tomb
+Of his self-love, to stop posterity?
+Thou art thy mother's glass, and she in thee
+Calls back the lovely April of her prime:
+So thou through windows of thine age shall see
+Despite of wrinkles this thy golden time.
+But if thou live, remember'd not to be,
+Die single, and thine image dies with thee.
+Unthrifty loveliness, why dost thou spend
+Upon thyself thy beauty's legacy?
+Nature's bequest gives nothing but doth lend,
+And being frank she lends to those are free.
+Then, beauteous niggard, why dost thou abuse
+The bounteous largess given thee to give?
+Profitless usurer, why dost thou use
+So great a sum of sums, yet canst not live?
+For having traffic with thyself alone,
+Thou of thyself thy sweet self dost deceive.
+Then how, when nature calls thee to be gone,
+What acceptable audit canst thou leave?
+Thy unused beauty must be tomb'd with thee,
+Which, used, lives th' executor to be.
+Those hours, that with gentle work did frame
+The lovely gaze where every eye doth dwell,
+Will play the tyrants to the very same
+And that unfair which fairly doth excel:
+For never-resting time leads summer on
+To hideous winter and confounds him there;
+Sap cheque'd with frost and lusty leaves quite gone,
+Beauty o'ersnow'd and bareness every where:
+Then, were not summer's distillation left,
+A liquid prisoner pent in walls of glass,
+Beauty's effect with beauty were bereft,
+Nor it nor no remembrance what it was:
+But flowers distill'd though they with winter meet,
+Leese but their show; their substance still lives sweet.
+Then let not winter's ragged hand deface
+In thee thy summer, ere thou be distill'd:
+Make sweet some vial; treasure thou some place
+With beauty's treasure, ere it be self-kill'd.
+That use is not forbidden usury,
+Which happies those that pay the willing loan;
+That's for thyself to breed another thee,
+Or ten times happier, be it ten for one;
+Ten times thyself were happier than thou art,
+If ten of thine ten times refigured thee:
+Then what could death do, if thou shouldst depart,
+Leaving thee living in posterity?
+Be not self-will'd, for thou art much too fair
+To be death's conquest and make worms thine heir.
+Lo! in the orient when the gracious light
+Lifts up his burning head, each under eye
+Doth homage to his new-appearing sight,
+Serving with looks his sacred majesty;
+And having climb'd the steep-up heavenly hill,
+Resembling strong youth in his middle age,
+yet mortal looks adore his beauty still,
+Attending on his golden pilgrimage;
+But when from highmost pitch, with weary car,
+Like feeble age, he reeleth from the day,
+The eyes, 'fore duteous, now converted are
+From his low tract and look another way:
+So thou, thyself out-going in thy noon,
+Unlook'd on diest, unless thou get a son.
+Music to hear, why hear'st thou music sadly?
+Sweets with sweets war not, joy delights in joy.
+Why lovest thou that which thou receivest not gladly,
+Or else receivest with pleasure thine annoy?
+If the true concord of well-tuned sounds,
+By unions married, do offend thine ear,
+They do but sweetly chide thee, who confounds
+In singleness the parts that thou shouldst bear.
+Mark how one string, sweet husband to another,
+Strikes each in each by mutual ordering,
+Resembling sire and child and happy mother
+Who all in one, one pleasing note do sing:
+Whose speechless song, being many, seeming one,
+Sings this to thee: 'thou single wilt prove none.'
+Is it for fear to wet a widow's eye
+That thou consumest thyself in single life?
+Ah! if thou issueless shalt hap to die.
+The world will wail thee, like a makeless wife;
+The world will be thy widow and still weep
+That thou no form of thee hast left behind,
+When every private widow well may keep
+By children's eyes her husband's shape in mind.
+Look, what an unthrift in the world doth spend
+Shifts but his place, for still the world enjoys it;
+But beauty's waste hath in the world an end,
+And kept unused, the user so destroys it.
+No love toward others in that bosom sits
+That on himself such murderous shame commits.
+For shame! deny that thou bear'st love to any,
+Who for thyself art so unprovident.
+Grant, if thou wilt, thou art beloved of many,
+But that thou none lovest is most evident;
+For thou art so possess'd with murderous hate
+That 'gainst thyself thou stick'st not to conspire.
+Seeking that beauteous roof to ruinate
+Which to repair should be thy chief desire.
+O, change thy thought, that I may change my mind!
+Shall hate be fairer lodged than gentle love?
+Be, as thy presence is, gracious and kind,
+Or to thyself at least kind-hearted prove:
+Make thee another self, for love of me,
+That beauty still may live in thine or thee.
+As fast as thou shalt wane, so fast thou growest
+In one of thine, from that which thou departest;
+And that fresh blood which youngly thou bestowest
+Thou mayst call thine when thou from youth convertest.
+Herein lives wisdom, beauty and increase:
+Without this, folly, age and cold decay:
+If all were minded so, the times should cease
+And threescore year would make the world away.
+Let those whom Nature hath not made for store,
+Harsh featureless and rude, barrenly perish:
+Look, whom she best endow'd she gave the more;
+Which bounteous gift thou shouldst in bounty cherish:
+She carved thee for her seal, and meant thereby
+Thou shouldst print more, not let that copy die.
+When I do count the clock that tells the time,
+And see the brave day sunk in hideous night;
+When I behold the violet past prime,
+And sable curls all silver'd o'er with white;
+When lofty trees I see barren of leaves
+Which erst from heat did canopy the herd,
+And summer's green all girded up in sheaves
+Borne on the bier with white and bristly beard,
+Then of thy beauty do I question make,
+That thou among the wastes of time must go,
+Since sweets and beauties do themselves forsake
+And die as fast as they see others grow;
+And nothing 'gainst Time's scythe can make defence
+Save breed, to brave him when he takes thee hence.
+O, that you were yourself! but, love, you are
+No longer yours than you yourself here live:
+Against this coming end you should prepare,
+And your sweet semblance to some other give.
+So should that beauty which you hold in lease
+Find no determination: then you were
+Yourself again after yourself's decease,
+When your sweet issue your sweet form should bear.
+Who lets so fair a house fall to decay,
+Which husbandry in honour might uphold
+Against the stormy gusts of winter's day
+And barren rage of death's eternal cold?
+O, none but unthrifts! Dear my love, you know
+You had a father: let your son say so.
+Not from the stars do I my judgment pluck;
+And yet methinks I have astronomy,
+But not to tell of good or evil luck,
+Of plagues, of dearths, or seasons' quality;
+Nor can I fortune to brief minutes tell,
+Pointing to each his thunder, rain and wind,
+Or say with princes if it shall go well,
+By oft predict that I in heaven find:
+But from thine eyes my knowledge I derive,
+And, constant stars, in them I read such art
+As truth and beauty shall together thrive,
+If from thyself to store thou wouldst convert;
+Or else of thee this I prognosticate:
+Thy end is truth's and beauty's doom and date.
+When I consider every thing that grows
+Holds in perfection but a little moment,
+That this huge stage presenteth nought but shows
+Whereon the stars in secret influence comment;
+When I perceive that men as plants increase,
+Cheered and cheque'd even by the self-same sky,
+Vaunt in their youthful sap, at height decrease,
+And wear their brave state out of memory;
+Then the conceit of this inconstant stay
+Sets you most rich in youth before my sight,
+Where wasteful Time debateth with Decay,
+To change your day of youth to sullied night;
+And all in war with Time for love of you,
+As he takes from you, I engraft you new.
+But wherefore do not you a mightier way
+Make war upon this bloody tyrant, Time?
+And fortify yourself in your decay
+With means more blessed than my barren rhyme?
+Now stand you on the top of happy hours,
+And many maiden gardens yet unset
+With virtuous wish would bear your living flowers,
+Much liker than your painted counterfeit:
+So should the lines of life that life repair,
+Which this, Time's pencil, or my pupil pen,
+Neither in inward worth nor outward fair,
+Can make you live yourself in eyes of men.
+To give away yourself keeps yourself still,
+And you must live, drawn by your own sweet skill.
+Who will believe my verse in time to come,
+If it were fill'd with your most high deserts?
+Though yet, heaven knows, it is but as a tomb
+Which hides your life and shows not half your parts.
+If I could write the beauty of your eyes
+And in fresh numbers number all your graces,
+The age to come would say 'This poet lies:
+Such heavenly touches ne'er touch'd earthly faces.'
+So should my papers yellow'd with their age
+Be scorn'd like old men of less truth than tongue,
+And your true rights be term'd a poet's rage
+And stretched metre of an antique song:
+But were some child of yours alive that time,
+You should live twice; in it and in my rhyme.
+Shall I compare thee to a summer's day?
+Thou art more lovely and more temperate:
+Rough winds do shake the darling buds of May,
+And summer's lease hath all too short a date:
+Sometime too hot the eye of heaven shines,
+And often is his gold complexion dimm'd;
+And every fair from fair sometime declines,
+By chance or nature's changing course untrimm'd;
+But thy eternal summer shall not fade
+Nor lose possession of that fair thou owest;
+Nor shall Death brag thou wander'st in his shade,
+When in eternal lines to time thou growest:
+So long as men can breathe or eyes can see,
+So long lives this and this gives life to thee.
+Devouring Time, blunt thou the lion's paws,
+And make the earth devour her own sweet brood;
+Pluck the keen teeth from the fierce tiger's jaws,
+And burn the long-lived phoenix in her blood;
+Make glad and sorry seasons as thou fleets,
+And do whate'er thou wilt, swift-footed Time,
+To the wide world and all her fading sweets;
+But I forbid thee one most heinous crime:
+O, carve not with thy hours my love's fair brow,
+Nor draw no lines there with thine antique pen;
+Him in thy course untainted do allow
+For beauty's pattern to succeeding men.
+Yet, do thy worst, old Time: despite thy wrong,
+My love shall in my verse ever live young.
+A woman's face with Nature's own hand painted
+Hast thou, the master-mistress of my passion;
+A woman's gentle heart, but not acquainted
+With shifting change, as is false women's fashion;
+An eye more bright than theirs, less false in rolling,
+Gilding the object whereupon it gazeth;
+A man in hue, all 'hues' in his controlling,
+Much steals men's eyes and women's souls amazeth.
+And for a woman wert thou first created;
+Till Nature, as she wrought thee, fell a-doting,
+And by addition me of thee defeated,
+By adding one thing to my purpose nothing.
+But since she prick'd thee out for women's pleasure,
+Mine be thy love and thy love's use their treasure.
+So is it not with me as with that Muse
+Stirr'd by a painted beauty to his verse,
+Who heaven itself for ornament doth use
+And every fair with his fair doth rehearse
+Making a couplement of proud compare,
+With sun and moon, with earth and sea's rich gems,
+With April's first-born flowers, and all things rare
+That heaven's air in this huge rondure hems.
+O' let me, true in love, but truly write,
+And then believe me, my love is as fair
+As any mother's child, though not so bright
+As those gold candles fix'd in heaven's air:
+Let them say more than like of hearsay well;
+I will not praise that purpose not to sell.
+My glass shall not persuade me I am old,
+So long as youth and thou are of one date;
+But when in thee time's furrows I behold,
+Then look I death my days should expiate.
+For all that beauty that doth cover thee
+Is but the seemly raiment of my heart,
+Which in thy breast doth live, as thine in me:
+How can I then be elder than thou art?
+O, therefore, love, be of thyself so wary
+As I, not for myself, but for thee will;
+Bearing thy heart, which I will keep so chary
+As tender nurse her babe from faring ill.
+Presume not on thy heart when mine is slain;
+Thou gavest me thine, not to give back again.
+As an unperfect actor on the stage
+Who with his fear is put besides his part,
+Or some fierce thing replete with too much rage,
+Whose strength's abundance weakens his own heart.
+So I, for fear of trust, forget to say
+The perfect ceremony of love's rite,
+And in mine own love's strength seem to decay,
+O'ercharged with burden of mine own love's might.
+O, let my books be then the eloquence
+And dumb presagers of my speaking breast,
+Who plead for love and look for recompense
+More than that tongue that more hath more express'd.
+O, learn to read what silent love hath writ:
+To hear with eyes belongs to love's fine wit.
+Mine eye hath play'd the painter and hath stell'd
+Thy beauty's form in table of my heart;
+My body is the frame wherein 'tis held,
+And perspective it is the painter's art.
+For through the painter must you see his skill,
+To find where your true image pictured lies;
+Which in my bosom's shop is hanging still,
+That hath his windows glazed with thine eyes.
+Now see what good turns eyes for eyes have done:
+Mine eyes have drawn thy shape, and thine for me
+Are windows to my breast, where-through the sun
+Delights to peep, to gaze therein on thee;
+Yet eyes this cunning want to grace their art;
+They draw but what they see, know not the heart.
+Let those who are in favour with their stars
+Of public honour and proud titles boast,
+Whilst I, whom fortune of such triumph bars,
+Unlook'd for joy in that I honour most.
+Great princes' favourites their fair leaves spread
+But as the marigold at the sun's eye,
+And in themselves their pride lies buried,
+For at a frown they in their glory die.
+The painful warrior famoused for fight,
+After a thousand victories once foil'd,
+Is from the book of honour razed quite,
+And all the rest forgot for which he toil'd:
+Then happy I, that love and am beloved
+Where I may not remove nor be removed.
+Lord of my love, to whom in vassalage
+Thy merit hath my duty strongly knit,
+To thee I send this written embassage,
+To witness duty, not to show my wit:
+Duty so great, which wit so poor as mine
+May make seem bare, in wanting words to show it,
+But that I hope some good conceit of thine
+In thy soul's thought, all naked, will bestow it;
+Till whatsoever star that guides my moving
+Points on me graciously with fair aspect
+And puts apparel on my tatter'd loving,
+To show me worthy of thy sweet respect:
+Then may I dare to boast how I do love thee;
+Till then not show my head where thou mayst prove me.
+Weary with toil, I haste me to my bed,
+The dear repose for limbs with travel tired;
+But then begins a journey in my head,
+To work my mind, when body's work's expired:
+For then my thoughts, from far where I abide,
+Intend a zealous pilgrimage to thee,
+And keep my drooping eyelids open wide,
+Looking on darkness which the blind do see
+Save that my soul's imaginary sight
+Presents thy shadow to my sightless view,
+Which, like a jewel hung in ghastly night,
+Makes black night beauteous and her old face new.
+Lo! thus, by day my limbs, by night my mind,
+For thee and for myself no quiet find.
+How can I then return in happy plight,
+That am debarr'd the benefit of rest?
+When day's oppression is not eased by night,
+But day by night, and night by day, oppress'd?
+And each, though enemies to either's reign,
+Do in consent shake hands to torture me;
+The one by toil, the other to complain
+How far I toil, still farther off from thee.
+I tell the day, to please them thou art bright
+And dost him grace when clouds do blot the heaven:
+So flatter I the swart-complexion'd night,
+When sparkling stars twire not thou gild'st the even.
+But day doth daily draw my sorrows longer
+And night doth nightly make grief's strength seem stronger.
+When, in disgrace with fortune and men's eyes,
+I all alone beweep my outcast state
+And trouble deal heaven with my bootless cries
+And look upon myself and curse my fate,
+Wishing me like to one more rich in hope,
+Featured like him, like him with friends possess'd,
+Desiring this man's art and that man's scope,
+With what I most enjoy contented least;
+Yet in these thoughts myself almost despising,
+Haply I think on thee, and then my state,
+Like to the lark at break of day arising
+From sullen earth, sings hymns at heaven's gate;
+For thy sweet love remember'd such wealth brings
+That then I scorn to change my state with kings.
+When to the sessions of sweet silent thought
+I summon up remembrance of things past,
+I sigh the lack of many a thing I sought,
+And with old woes new wail my dear time's waste:
+Then can I drown an eye, unused to flow,
+For precious friends hid in death's dateless night,
+And weep afresh love's long since cancell'd woe,
+And moan the expense of many a vanish'd sight:
+Then can I grieve at grievances foregone,
+And heavily from woe to woe tell o'er
+The sad account of fore-bemoaned moan,
+Which I new pay as if not paid before.
+But if the while I think on thee, dear friend,
+All losses are restored and sorrows end.
+Thy bosom is endeared with all hearts,
+Which I by lacking have supposed dead,
+And there reigns love and all love's loving parts,
+And all those friends which I thought buried.
+How many a holy and obsequious tear
+Hath dear religious love stol'n from mine eye
+As interest of the dead, which now appear
+But things removed that hidden in thee lie!
+Thou art the grave where buried love doth live,
+Hung with the trophies of my lovers gone,
+Who all their parts of me to thee did give;
+That due of many now is thine alone:
+Their images I loved I view in thee,
+And thou, all they, hast all the all of me.
+If thou survive my well-contented day,
+When that churl Death my bones with dust shall cover,
+And shalt by fortune once more re-survey
+These poor rude lines of thy deceased lover,
+Compare them with the bettering of the time,
+And though they be outstripp'd by every pen,
+Reserve them for my love, not for their rhyme,
+Exceeded by the height of happier men.
+O, then vouchsafe me but this loving thought:
+'Had my friend's Muse grown with this growing age,
+A dearer birth than this his love had brought,
+To march in ranks of better equipage:
+But since he died and poets better prove,
+Theirs for their style I'll read, his for his love.'
+Full many a glorious morning have I seen
+Flatter the mountain-tops with sovereign eye,
+Kissing with golden face the meadows green,
+Gilding pale streams with heavenly alchemy;
+Anon permit the basest clouds to ride
+With ugly rack on his celestial face,
+And from the forlorn world his visage hide,
+Stealing unseen to west with this disgrace:
+Even so my sun one early morn did shine
+With all triumphant splendor on my brow;
+But out, alack! he was but one hour mine;
+The region cloud hath mask'd him from me now.
+Yet him for this my love no whit disdaineth;
+Suns of the world may stain when heaven's sun staineth.
+Why didst thou promise such a beauteous day,
+And make me travel forth without my cloak,
+To let base clouds o'ertake me in my way,
+Hiding thy bravery in their rotten smoke?
+'Tis not enough that through the cloud thou break,
+To dry the rain on my storm-beaten face,
+For no man well of such a salve can speak
+That heals the wound and cures not the disgrace:
+Nor can thy shame give physic to my grief;
+Though thou repent, yet I have still the loss:
+The offender's sorrow lends but weak relief
+To him that bears the strong offence's cross.
+Ah! but those tears are pearl which thy love sheds,
+And they are rich and ransom all ill deeds.
+No more be grieved at that which thou hast done:
+Roses have thorns, and silver fountains mud;
+Clouds and eclipses stain both moon and sun,
+And loathsome canker lives in sweetest bud.
+All men make faults, and even I in this,
+Authorizing thy trespass with compare,
+Myself corrupting, salving thy amiss,
+Excusing thy sins more than thy sins are;
+For to thy sensual fault I bring in sense--
+Thy adverse party is thy advocate--
+And 'gainst myself a lawful plea commence:
+Such civil war is in my love and hate
+That I an accessary needs must be
+To that sweet thief which sourly robs from me.
+Let me confess that we two must be twain,
+Although our undivided loves are one:
+So shall those blots that do with me remain
+Without thy help by me be borne alone.
+In our two loves there is but one respect,
+Though in our lives a separable spite,
+Which though it alter not love's sole effect,
+Yet doth it steal sweet hours from love's delight.
+I may not evermore acknowledge thee,
+Lest my bewailed guilt should do thee shame,
+Nor thou with public kindness honour me,
+Unless thou take that honour from thy name:
+But do not so; I love thee in such sort
+As, thou being mine, mine is thy good report.
+As a decrepit father takes delight
+To see his active child do deeds of youth,
+So I, made lame by fortune's dearest spite,
+Take all my comfort of thy worth and truth.
+For whether beauty, birth, or wealth, or wit,
+Or any of these all, or all, or more,
+Entitled in thy parts do crowned sit,
+I make my love engrafted to this store:
+So then I am not lame, poor, nor despised,
+Whilst that this shadow doth such substance give
+That I in thy abundance am sufficed
+And by a part of all thy glory live.
+Look, what is best, that best I wish in thee:
+This wish I have; then ten times happy me!
@@ -0,0 +1,303 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+use std::sync::Arc;
+
+use rand::rngs::StdRng;
+use rand::seq::SliceRandom;
+use rand::{Rng, SeedableRng};
+
+use super::SampleRequest;
+use crate::cli::SpeedBenchConfig;
+use crate::error::{BenchError, Result};
+use crate::tokenizer::TokenizerKind;
+
+/// Marker text for masked entries that need external fetch.
+const MASKED_PREFIX: &str = "FULL BENCHMARK DATA SHOULD BE FETCHED";
+
+/// Cache directory for downloaded SPEED-Bench datasets.
+fn cache_dir() -> std::path::PathBuf {
+    dirs::cache_dir()
+        .unwrap_or_else(|| std::path::PathBuf::from("/tmp"))
+        .join("vllm-bench")
+        .join("datasets")
+}
+
+/// Download SPEED-Bench dataset from HuggingFace datasets-server API.
+/// Results are cached as JSON locally for subsequent runs.
+pub fn download_speed_bench(config: SpeedBenchConfig) -> Result<String> {
+    let config_name = config.as_str();
+
+    let dir = cache_dir();
+    std::fs::create_dir_all(&dir)?;
+    let cache_path = dir.join(format!("speed-bench-{config_name}.json"));
+
+    // Return cached file if it exists
+    if cache_path.exists() {
+        let path_str = cache_path.to_string_lossy().to_string();
+        println!("SPEED-Bench ({config_name}) cached: {path_str}");
+        return Ok(path_str);
+    }
+
+    println!("Downloading SPEED-Bench ({config_name}) from HuggingFace datasets-server...");
+
+    let client = reqwest::blocking::Client::builder()
+        .timeout(std::time::Duration::from_secs(120))
+        .build()
+        .map_err(|e| BenchError::Config(format!("Failed to build HTTP client: {e}")))?;
+
+    let mut all_rows: Vec<serde_json::Value> = Vec::new();
+    let mut offset = 0usize;
+    let page_size = 100usize;
+
+    loop {
+        let url = format!(
+            "https://datasets-server.huggingface.co/rows\
+             ?dataset=nvidia/SPEED-Bench\
+             &config={config_name}\
+             &split=test\
+             &offset={offset}\
+             &length={page_size}"
+        );
+
+        // Retry on transient errors (502, 503, timeouts)
+        let max_retries = 3;
+        let mut data: Option<serde_json::Value> = None;
+        for attempt in 0..=max_retries {
+            let resp = match client.get(&url).send() {
+                Ok(r) => r,
+                Err(e) => {
+                    if attempt < max_retries {
+                        std::thread::sleep(std::time::Duration::from_secs(
+                            2 * (attempt as u64 + 1),
+                        ));
+                        continue;
+                    }
+                    return Err(BenchError::Config(format!(
+                        "SPEED-Bench download failed after {max_retries} retries: {e}"
+                    )));
+                }
+            };
+
+            if resp.status().is_server_error() && attempt < max_retries {
+                std::thread::sleep(std::time::Duration::from_secs(2 * (attempt as u64 + 1)));
+                continue;
+            }
+
+            if !resp.status().is_success() {
+                return Err(BenchError::Config(format!(
+                    "SPEED-Bench API returned HTTP {}",
+                    resp.status()
+                )));
+            }
+
+            data = Some(resp.json().map_err(|e| {
+                BenchError::Config(format!("Failed to parse SPEED-Bench API response: {e}"))
+            })?);
+            break;
+        }
+
+        let data = data.unwrap();
+
+        let rows = data["rows"]
+            .as_array()
+            .ok_or_else(|| BenchError::Config("No 'rows' in API response".into()))?;
+
+        if rows.is_empty() {
+            break;
+        }
+
+        for row in rows {
+            if let Some(row_data) = row.get("row") {
+                all_rows.push(row_data.clone());
+            }
+        }
+
+        let fetched = rows.len();
+        offset += fetched;
+
+        // Print progress
+        let total = data["num_rows_total"].as_u64().unwrap_or(0);
+        eprint!("\r  Fetched {offset}/{total} rows...");
+
+        if fetched < page_size {
+            break;
+        }
+    }
+    eprintln!(); // newline after progress
+
+    if all_rows.is_empty() {
+        return Err(BenchError::Config(
+            "SPEED-Bench download returned no rows".into(),
+        ));
+    }
+
+    // Save to cache
+    let json_str = serde_json::to_string(&all_rows)?;
+    std::fs::write(&cache_path, &json_str)?;
+
+    let path_str = cache_path.to_string_lossy().to_string();
+    println!(
+        "SPEED-Bench ({config_name}): {} rows saved to {path_str}",
+        all_rows.len()
+    );
+    Ok(path_str)
+}
+
+/// Load SPEED-Bench dataset and convert to SampleRequests.
+///
+/// Filters out masked entries and optionally filters by category.
+/// Requires an output length override since SPEED-Bench has no reference outputs.
+pub fn load_speed_bench_dataset(
+    tokenizer: &TokenizerKind,
+    dataset_path: &str,
+    num_requests: usize,
+    output_len: usize,
+    seed: u64,
+    request_id_prefix: &str,
+    category_filter: Option<&str>,
+    no_oversample: bool,
+    disable_shuffle: bool,
+    max_input_len: Option<usize>,
+) -> Result<Vec<SampleRequest>> {
+    let content = std::fs::read_to_string(dataset_path).map_err(|e| {
+        BenchError::Config(format!(
+            "Failed to read SPEED-Bench file '{dataset_path}': {e}"
+        ))
+    })?;
+
+    let entries: Vec<serde_json::Value> = serde_json::from_str(&content)
+        .map_err(|e| BenchError::Config(format!("Invalid JSON in SPEED-Bench file: {e}")))?;
+
+    // Filter entries
+    let mut filtered: Vec<&serde_json::Value> = entries
+        .iter()
+        .filter(|entry| {
+            // Must have turns array with at least one non-empty entry
+            let turns = match entry.get("turns").and_then(|t| t.as_array()) {
+                Some(t) if !t.is_empty() => t,
+                _ => return false,
+            };
+
+            // Skip masked entries
+            let first_turn = turns[0].as_str().unwrap_or("");
+            if first_turn.starts_with(MASKED_PREFIX) || first_turn.is_empty() {
+                return false;
+            }
+
+            // Single-turn only: skip multi-turn entries
+            let is_multiturn = entry.get("multiturn").and_then(|m| m.as_bool()).unwrap_or(false);
+            if is_multiturn {
+                return false;
+            }
+
+            // Category filter
+            if let Some(cat) = category_filter {
+                let entry_cat = entry.get("category").and_then(|c| c.as_str()).unwrap_or("");
+                if entry_cat != cat {
+                    return false;
+                }
+            }
+
+            true
+        })
+        .collect();
+
+    if filtered.is_empty() {
+        let cat_msg = category_filter.map(|c| format!(" with category '{c}'")).unwrap_or_default();
+        return Err(BenchError::Config(format!(
+            "No valid single-turn entries in SPEED-Bench{cat_msg}. \
+             Try a different --speed-bench-config or remove --speed-bench-category filter."
+        )));
+    }
+
+    // Shuffle
+    let mut rng = StdRng::seed_from_u64(seed);
+    if !disable_shuffle {
+        filtered.shuffle(&mut rng);
+    }
+
+    // Build SampleRequests
+    let mut samples = Vec::new();
+    let mut idx = 0;
+
+    for entry in &filtered {
+        if samples.len() >= num_requests {
+            break;
+        }
+
+        let turns = entry["turns"].as_array().unwrap();
+        let prompt = turns[0].as_str().unwrap_or("");
+
+        // Tokenize to get prompt length
+        let prompt_ids = tokenizer.encode(prompt, false)?;
+        let prompt_len = prompt_ids.len();
+
+        if prompt_len < 4 {
+            continue;
+        }
+
+        // Truncate if max_input_len is set
+        let (final_prompt, final_len) = if let Some(max_len) = max_input_len {
+            if prompt_len > max_len {
+                let truncated_ids = &prompt_ids[..max_len];
+                let truncated_text = tokenizer.decode(truncated_ids, true)?;
+                (Arc::from(truncated_text.as_str()), max_len)
+            } else {
+                (Arc::from(prompt), prompt_len)
+            }
+        } else {
+            (Arc::from(prompt), prompt_len)
+        };
+
+        samples.push(SampleRequest {
+            prompt: final_prompt,
+            prompt_len: final_len,
+            expected_output_len: output_len,
+            request_id: Some(format!("{request_id_prefix}{idx}")),
+            ..Default::default()
+        });
+        idx += 1;
+    }
+
+    // Oversample if needed
+    if samples.len() < num_requests {
+        if no_oversample {
+            println!(
+                "Skipping oversampling. Total samples: {} (requested: {num_requests})",
+                samples.len()
+            );
+        } else if !samples.is_empty() {
+            let original_len = samples.len();
+            let needed = num_requests - original_len;
+            for i in 0..needed {
+                let mut req = samples[rng.random_range(0..original_len)].clone();
+                req.request_id = Some(format!("{request_id_prefix}{}", original_len + i));
+                samples.push(req);
+            }
+            println!(
+                "Oversampled SPEED-Bench from {original_len} to {} total samples.",
+                samples.len()
+            );
+        }
+    }
+
+    if samples.is_empty() {
+        return Err(BenchError::Config(
+            "No valid samples after filtering SPEED-Bench dataset.".into(),
+        ));
+    }
+
+    // Print category distribution
+    let mut cat_counts: std::collections::HashMap<&str, usize> = std::collections::HashMap::new();
+    for entry in &filtered[..filtered.len().min(samples.len())] {
+        let cat = entry.get("category").and_then(|c| c.as_str()).unwrap_or("unknown");
+        *cat_counts.entry(cat).or_insert(0) += 1;
+    }
+    let mut cats: Vec<_> = cat_counts.into_iter().collect();
+    cats.sort_by_key(|b| std::cmp::Reverse(b.1));
+    let cat_str: Vec<String> = cats.iter().map(|(k, v)| format!("{k}:{v}")).collect();
+    println!("SPEED-Bench categories: {}", cat_str.join(", "));
+
+    Ok(samples)
+}
@@ -0,0 +1,37 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+use thiserror::Error;
+
+#[derive(Error, Debug)]
+pub enum BenchError {
+    #[error("HTTP request failed: {0}")]
+    Http(#[from] reqwest::Error),
+
+    #[error("JSON error: {0}")]
+    Json(#[from] serde_json::Error),
+
+    #[error("Tokenizer error: {0}")]
+    Tokenizer(String),
+
+    /// The server's /tokenize//detokenize endpoint is not usable (4xx status:
+    /// not exposed, or rejected by a gateway such as LLM-d/EPP that returns
+    /// 400 instead of 404). Callers treat this as "skip verification", unlike
+    /// `Tokenizer` errors which are genuine failures.
+    #[error("tokenize endpoint unavailable: {0}")]
+    TokenizeUnavailable(String),
+
+    #[error("Configuration error: {0}")]
+    Config(String),
+
+    #[error("Endpoint not ready after {0}s: {1}")]
+    EndpointTimeout(u64, String),
+
+    #[error("Backend error: {0}")]
+    Backend(String),
+
+    #[error("IO error: {0}")]
+    Io(#[from] std::io::Error),
+}
+
+pub type Result<T> = std::result::Result<T, BenchError>;
@@ -0,0 +1,54 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+//! Sync facade over the async `hf_hub` API.
+//!
+//! The workspace bans rustls (`rust/deny.toml`), but hf-hub's sync `ureq`
+//! backend unconditionally pulls ureq's default rustls feature. So we use the
+//! reqwest/native-tls tokio API instead, and bridge blocking callers (dataset
+//! loaders, tokenizer fallback in rayon threads) by running each download on a
+//! dedicated thread with its own single-threaded runtime.
+
+use std::path::PathBuf;
+
+/// A handle to a HuggingFace Hub repo, downloading via hf-hub's on-disk cache.
+pub struct HubRepo {
+    repo: hf_hub::Repo,
+}
+
+impl HubRepo {
+    pub fn model(model_id: String) -> Self {
+        Self {
+            repo: hf_hub::Repo::model(model_id),
+        }
+    }
+
+    pub fn dataset(repo_id: String) -> Self {
+        Self {
+            repo: hf_hub::Repo::dataset(repo_id),
+        }
+    }
+
+    /// Download (or fetch from cache) a single file from the repo.
+    /// Auth is handled by hf-hub via HF_TOKEN / the cached login token.
+    pub fn get(&self, filename: &str) -> Result<PathBuf, String> {
+        let repo = self.repo.clone();
+        let filename = filename.to_string();
+        std::thread::spawn(move || {
+            let rt = tokio::runtime::Builder::new_current_thread()
+                .enable_all()
+                .build()
+                .map_err(|e| format!("Failed to build download runtime: {e}"))?;
+            rt.block_on(async move {
+                let mut builder = hf_hub::api::tokio::ApiBuilder::from_env();
+                if let Ok(token) = std::env::var("HF_TOKEN") {
+                    builder = builder.with_token(Some(token));
+                }
+                let api = builder.build().map_err(|e| format!("Failed to init HF API: {e}"))?;
+                api.repo(repo).get(&filename).await.map_err(|e| format!("{e}"))
+            })
+        })
+        .join()
+        .map_err(|_| "HF Hub download thread panicked".to_string())?
+    }
+}
@@ -0,0 +1,86 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+mod backends;
+mod benchmark;
+mod cli;
+mod compare;
+mod config;
+mod datasets;
+mod error;
+mod hub;
+mod metrics;
+mod multi_run;
+mod multi_turn;
+mod output;
+mod rate_control;
+mod ready_checker;
+mod sweep;
+mod tiktoken;
+mod tokenizer;
+
+use anyhow::Context;
+
+pub use cli::{
+    BackendKind, BenchServeArgs, DatasetName, LoraAssignment, RampUpStrategy, SpeedBenchConfig,
+};
+use config::BenchConfig;
+
+/// Prepare process-wide resources for a benchmark run.
+pub fn prepare_process() {
+    // Raise the open-file soft limit to the hard limit. High-concurrency
+    // benchmarks (1024+ requests) easily exceed the default 1024 fd soft limit.
+    if let Ok(new) = rlimit::increase_nofile_limit(u64::MAX)
+        && new > 1024
+    {
+        eprintln!("Open-file limit: {new}");
+    }
+}
+
+/// Run the online serving benchmark.
+pub async fn run(args: BenchServeArgs) -> anyhow::Result<()> {
+    // --- Compare mode: no server needed, just diff two JSON files ---
+    if let Some(ref files) = args.compare {
+        return compare::compare_results(&files[0], &files[1]).context("Comparison failed");
+    }
+
+    let config = BenchConfig::from_args(&args).context("Configuration error")?;
+
+    async {
+        if config.multi_turn {
+            if let Some(ref sweep_mc) = args.sweep_max_concurrency {
+                // --- Sweep over concurrency in multi-turn mode ---
+                let values = sweep::parse_concurrency_values(sweep_mc)
+                    .context("Invalid --sweep-max-concurrency")?;
+                sweep::run_multi_turn_concurrency_sweep(
+                    &config,
+                    &values,
+                    args.sweep_num_prompts_factor,
+                )
+                .await?;
+            } else {
+                // --- Single multi-turn conversation benchmark ---
+                multi_turn::run_multi_turn_benchmark(&config).await?;
+            }
+        } else if let Some(ref sweep_mc) = args.sweep_max_concurrency {
+            // --- Sweep over max-concurrency ---
+            let values = sweep::parse_concurrency_values(sweep_mc)
+                .context("Invalid --sweep-max-concurrency")?;
+            sweep::run_concurrency_sweep(&config, &values, args.sweep_num_prompts_factor).await?;
+        } else if let Some(ref sweep_rate) = args.sweep_request_rate {
+            // --- Sweep over request-rate ---
+            let values =
+                sweep::parse_rate_values(sweep_rate).context("Invalid --sweep-request-rate")?;
+            sweep::run_rate_sweep(&config, &values).await?;
+        } else if args.num_runs > 1 {
+            // --- Multi-run with statistical aggregation ---
+            multi_run::run_multi(&config, args.num_runs).await?;
+        } else {
+            // --- Normal single benchmark ---
+            benchmark::run_benchmark(&config).await?;
+        }
+        anyhow::Ok(())
+    }
+    .await
+    .context("Benchmark failed")
+}
@@ -0,0 +1,32 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+#[cfg(not(target_env = "msvc"))]
+#[global_allocator]
+static GLOBAL: mimalloc::MiMalloc = mimalloc::MiMalloc;
+
+use anyhow::Context;
+use clap::Parser;
+
+#[derive(Parser)]
+#[command(
+    name = "vllm-bench",
+    about = "Benchmark online serving throughput",
+    version
+)]
+struct Cli {
+    #[command(flatten)]
+    args: vllm_bench::BenchServeArgs,
+}
+
+fn main() -> anyhow::Result<()> {
+    let cli = Cli::parse();
+    vllm_bench::prepare_process();
+
+    let runtime = tokio::runtime::Builder::new_multi_thread()
+        .enable_all()
+        .build()
+        .context("Failed to build tokio runtime")?;
+
+    runtime.block_on(vllm_bench::run(cli.args))
+}
@@ -0,0 +1,503 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+use crate::backends::RequestFuncOutput;
+use crate::config::GoodputConfig;
+use crate::datasets::SampleRequest;
+use crate::metrics::{BenchmarkMetrics, MultiTurnMetrics};
+use crate::multi_turn::ConversationOutput;
+
+/// Calculate benchmark metrics from request outputs.
+///
+/// Mirrors Python's `calculate_metrics()` from serve.py:392-599.
+pub fn calculate_metrics(
+    input_requests: &[SampleRequest],
+    outputs: &[RequestFuncOutput],
+    dur_s: f64,
+    selected_percentiles: &[f64],
+    _has_tokenizer: bool,
+    goodput_config: &GoodputConfig,
+) -> (BenchmarkMetrics, Vec<usize>) {
+    let mut actual_output_lens: Vec<usize> = Vec::with_capacity(outputs.len());
+    let mut total_input: usize = 0;
+    let mut completed: usize = 0;
+    let mut itls: Vec<f64> = Vec::new();
+    let mut tpots: Vec<f64> = Vec::new();
+    // Per-request TPOT for goodput SLO checking (parallel to ttfts/e2els).
+    let mut all_tpots: Vec<f64> = Vec::new();
+    let mut ttfts: Vec<f64> = Vec::new();
+    let mut e2els: Vec<f64> = Vec::new();
+
+    for (i, output) in outputs.iter().enumerate() {
+        if output.success {
+            let output_len = if output.output_tokens > 0 {
+                output.output_tokens
+            } else {
+                // Fallback: first token + ITL entries.
+                // Python re-encodes generated_text when tokenizer is available,
+                // but with vLLM's stream_options.include_usage=true,
+                // output_tokens is always set so this path is rarely hit.
+                1 + output.itl.len()
+            };
+
+            actual_output_lens.push(output_len);
+            total_input += input_requests[i].prompt_len;
+
+            if output_len > 1 {
+                let latency_minus_ttft = output.latency - output.ttft;
+                let tpot = latency_minus_ttft / (output_len as f64 - 1.0);
+                tpots.push(tpot);
+                all_tpots.push(tpot);
+            } else {
+                all_tpots.push(0.0);
+            }
+
+            itls.extend_from_slice(&output.itl);
+            ttfts.push(output.ttft);
+            e2els.push(output.latency);
+            completed += 1;
+        } else {
+            actual_output_lens.push(0);
+        }
+    }
+
+    let failed = outputs.len() - completed;
+
+    // Print failed request errors (capped to 10)
+    let failed_outputs: Vec<&RequestFuncOutput> = outputs.iter().filter(|o| !o.success).collect();
+    if !failed_outputs.is_empty() {
+        eprintln!("Failed requests during benchmark run detected (capping to 10):");
+        for (i, err) in failed_outputs.iter().take(10).enumerate() {
+            eprintln!("Error {i}: {}", err.error);
+        }
+    }
+
+    // Calculate max output tokens per second and max concurrent requests
+    let mut max_output_tokens_per_s = 0.0_f64;
+    let mut max_concurrent_requests: usize = 0;
+
+    let successful_outputs: Vec<&RequestFuncOutput> =
+        outputs.iter().filter(|o| o.success).collect();
+
+    if !successful_outputs.is_empty() {
+        let min_start_time =
+            successful_outputs.iter().map(|o| o.start_time).fold(f64::INFINITY, f64::min);
+        let max_end_time = successful_outputs
+            .iter()
+            .map(|o| o.start_time + o.latency)
+            .fold(f64::NEG_INFINITY, f64::max);
+
+        let raw_duration = (max_end_time - min_start_time).ceil() as usize + 1;
+        // Cap at 24 hours to prevent OOM from corrupted timing data
+        let duration_seconds = raw_duration.min(86_400);
+        let mut tokens_per_second = vec![0.0_f64; duration_seconds];
+        let mut concurrent_per_second = vec![0usize; duration_seconds];
+
+        for output in &successful_outputs {
+            // Calculate token generation timestamps
+            let mut token_times = vec![output.start_time + output.ttft];
+            let mut current_time = token_times[0];
+            for itl_value in &output.itl {
+                current_time += itl_value;
+                token_times.push(current_time);
+            }
+
+            // Add tokens to second buckets
+            for token_time in &token_times {
+                let bucket = (token_time - min_start_time) as usize;
+                if bucket < duration_seconds {
+                    tokens_per_second[bucket] += 1.0;
+                }
+            }
+
+            // Track concurrent requests
+            let start_second = (output.start_time - min_start_time) as usize;
+            let end_second = ((output.start_time + output.latency) - min_start_time) as usize;
+            for slot in concurrent_per_second
+                .iter_mut()
+                .take(end_second.min(duration_seconds - 1) + 1)
+                .skip(start_second)
+            {
+                *slot += 1;
+            }
+        }
+
+        max_output_tokens_per_s = tokens_per_second.iter().cloned().fold(0.0_f64, f64::max);
+        max_concurrent_requests = *concurrent_per_second.iter().max().unwrap_or(&0);
+    }
+
+    let total_output: usize = actual_output_lens.iter().sum();
+
+    // Compute goodput: count requests meeting ALL specified SLOs.
+    // Mirrors Python serve.py:458-481.
+    let good_completed = if !goodput_config.is_empty() {
+        let mut good = 0usize;
+        // ttfts, all_tpots, e2els are parallel (one per successful request)
+        for i in 0..ttfts.len() {
+            let mut is_good = true;
+            if let Some(slo_ms) = goodput_config.ttft_ms
+                && ttfts[i] * 1000.0 > slo_ms
+            {
+                is_good = false;
+            }
+            if let Some(slo_ms) = goodput_config.tpot_ms
+                && all_tpots[i] * 1000.0 > slo_ms
+            {
+                is_good = false;
+            }
+            if let Some(slo_ms) = goodput_config.e2el_ms
+                && e2els[i] * 1000.0 > slo_ms
+            {
+                is_good = false;
+            }
+            if is_good {
+                good += 1;
+            }
+        }
+        good
+    } else {
+        0
+    };
+
+    // Sort each metric array once, then compute median + percentiles from sorted data
+    let sorted_ttfts = sort_clone(&ttfts);
+    let sorted_tpots = sort_clone(&tpots);
+    let sorted_itls = sort_clone(&itls);
+    let sorted_e2els = sort_clone(&e2els);
+
+    let request_goodput = if !goodput_config.is_empty() {
+        good_completed as f64 / dur_s
+    } else {
+        0.0
+    };
+
+    let metrics = BenchmarkMetrics {
+        completed,
+        failed,
+        total_input,
+        total_output,
+        request_throughput: completed as f64 / dur_s,
+        request_goodput,
+        input_throughput: total_input as f64 / dur_s,
+        output_throughput: total_output as f64 / dur_s,
+        total_token_throughput: (total_input + total_output) as f64 / dur_s,
+        mean_ttft_ms: mean(&ttfts) * 1000.0,
+        median_ttft_ms: median_sorted(&sorted_ttfts) * 1000.0,
+        std_ttft_ms: std_dev(&ttfts) * 1000.0,
+        percentiles_ttft_ms: percentiles_from_sorted(&sorted_ttfts, selected_percentiles),
+        mean_tpot_ms: mean(&tpots) * 1000.0,
+        median_tpot_ms: median_sorted(&sorted_tpots) * 1000.0,
+        std_tpot_ms: std_dev(&tpots) * 1000.0,
+        percentiles_tpot_ms: percentiles_from_sorted(&sorted_tpots, selected_percentiles),
+        mean_itl_ms: mean(&itls) * 1000.0,
+        median_itl_ms: median_sorted(&sorted_itls) * 1000.0,
+        std_itl_ms: std_dev(&itls) * 1000.0,
+        percentiles_itl_ms: percentiles_from_sorted(&sorted_itls, selected_percentiles),
+        mean_e2el_ms: mean(&e2els) * 1000.0,
+        median_e2el_ms: median_sorted(&sorted_e2els) * 1000.0,
+        std_e2el_ms: std_dev(&e2els) * 1000.0,
+        percentiles_e2el_ms: percentiles_from_sorted(&sorted_e2els, selected_percentiles),
+        max_output_tokens_per_s,
+        max_concurrent_requests,
+        steady_state: None,
+    };
+
+    (metrics, actual_output_lens)
+}
+
+fn mean(data: &[f64]) -> f64 {
+    if data.is_empty() {
+        return 0.0;
+    }
+    data.iter().sum::<f64>() / data.len() as f64
+}
+
+fn sort_clone(data: &[f64]) -> Vec<f64> {
+    let mut sorted = data.to_vec();
+    sorted.sort_by(|a, b| a.partial_cmp(b).unwrap_or(std::cmp::Ordering::Equal));
+    sorted
+}
+
+fn median_sorted(sorted: &[f64]) -> f64 {
+    if sorted.is_empty() {
+        return 0.0;
+    }
+    let mid = sorted.len() / 2;
+    if sorted.len().is_multiple_of(2) {
+        (sorted[mid - 1] + sorted[mid]) / 2.0
+    } else {
+        sorted[mid]
+    }
+}
+
+fn std_dev(data: &[f64]) -> f64 {
+    if data.is_empty() {
+        return 0.0;
+    }
+    let m = mean(data);
+    let variance = data.iter().map(|x| (x - m).powi(2)).sum::<f64>() / data.len() as f64;
+    variance.sqrt()
+}
+
+fn percentile_sorted(sorted_data: &[f64], p: f64) -> f64 {
+    if sorted_data.is_empty() {
+        return 0.0;
+    }
+    if sorted_data.len() == 1 {
+        return sorted_data[0];
+    }
+    // Use the same interpolation as numpy (linear)
+    let idx = p / 100.0 * (sorted_data.len() - 1) as f64;
+    let lo = idx.floor() as usize;
+    let hi = idx.ceil() as usize;
+    let frac = idx - lo as f64;
+    if lo == hi {
+        sorted_data[lo]
+    } else {
+        sorted_data[lo] * (1.0 - frac) + sorted_data[hi] * frac
+    }
+}
+
+fn percentiles_from_sorted(sorted: &[f64], percentiles: &[f64]) -> Vec<(f64, f64)> {
+    if sorted.is_empty() {
+        return percentiles.iter().map(|&p| (p, 0.0)).collect();
+    }
+    percentiles
+        .iter()
+        .map(|&p| (p, percentile_sorted(sorted, p) * 1000.0))
+        .collect()
+}
+
+/// Calculate benchmark metrics for embedding/pooling requests.
+///
+/// Mirrors Python's `calculate_metrics_for_embeddings()` from serve.py.
+/// Key differences from `calculate_metrics()`:
+/// - Uses `output.prompt_len` (server-reported) instead of `input_requests[i].prompt_len`
+/// - Only computes E2EL (no TTFT/TPOT/ITL since pooling is non-streaming)
+/// - `total_output` is 0 (pooling produces no output tokens)
+/// - `total_token_throughput` = `total_input / dur_s` (input tokens only)
+pub fn calculate_embedding_metrics(
+    outputs: &[RequestFuncOutput],
+    dur_s: f64,
+    selected_percentiles: &[f64],
+) -> BenchmarkMetrics {
+    let mut total_input: usize = 0;
+    let mut completed: usize = 0;
+    let mut e2els: Vec<f64> = Vec::new();
+
+    for output in outputs {
+        if output.success {
+            e2els.push(output.latency);
+            completed += 1;
+            total_input += output.prompt_len;
+        }
+    }
+
+    let failed = outputs.len() - completed;
+
+    // Print failed request errors (capped to 10)
+    let failed_outputs: Vec<&RequestFuncOutput> = outputs.iter().filter(|o| !o.success).collect();
+    if !failed_outputs.is_empty() {
+        eprintln!("Failed requests during benchmark run detected (capping to 10):");
+        for (i, err) in failed_outputs.iter().take(10).enumerate() {
+            eprintln!("Error {i}: {}", err.error);
+        }
+    }
+
+    // Compute peak concurrent requests from start_time + latency windows
+    let successful_outputs: Vec<&RequestFuncOutput> =
+        outputs.iter().filter(|o| o.success).collect();
+    let max_concurrent_requests = if !successful_outputs.is_empty() {
+        let min_start =
+            successful_outputs.iter().map(|o| o.start_time).fold(f64::INFINITY, f64::min);
+        let max_end = successful_outputs
+            .iter()
+            .map(|o| o.start_time + o.latency)
+            .fold(f64::NEG_INFINITY, f64::max);
+
+        let raw_duration = (max_end - min_start).ceil() as usize + 1;
+        let duration_seconds = raw_duration.min(86_400);
+        let mut concurrent_per_second = vec![0usize; duration_seconds];
+
+        for output in &successful_outputs {
+            let start_second = (output.start_time - min_start) as usize;
+            let end_second = ((output.start_time + output.latency) - min_start) as usize;
+            for slot in concurrent_per_second
+                .iter_mut()
+                .take(end_second.min(duration_seconds - 1) + 1)
+                .skip(start_second)
+            {
+                *slot += 1;
+            }
+        }
+
+        *concurrent_per_second.iter().max().unwrap_or(&0)
+    } else {
+        0
+    };
+
+    let sorted_e2els = sort_clone(&e2els);
+
+    BenchmarkMetrics {
+        completed,
+        failed,
+        total_input,
+        total_output: 0,
+        request_throughput: completed as f64 / dur_s,
+        request_goodput: 0.0,
+        input_throughput: total_input as f64 / dur_s,
+        output_throughput: 0.0,
+        total_token_throughput: total_input as f64 / dur_s,
+        mean_ttft_ms: 0.0,
+        median_ttft_ms: 0.0,
+        std_ttft_ms: 0.0,
+        percentiles_ttft_ms: Vec::new(),
+        mean_tpot_ms: 0.0,
+        median_tpot_ms: 0.0,
+        std_tpot_ms: 0.0,
+        percentiles_tpot_ms: Vec::new(),
+        mean_itl_ms: 0.0,
+        median_itl_ms: 0.0,
+        std_itl_ms: 0.0,
+        percentiles_itl_ms: Vec::new(),
+        mean_e2el_ms: mean(&e2els) * 1000.0,
+        median_e2el_ms: median_sorted(&sorted_e2els) * 1000.0,
+        std_e2el_ms: std_dev(&e2els) * 1000.0,
+        percentiles_e2el_ms: percentiles_from_sorted(&sorted_e2els, selected_percentiles),
+        max_output_tokens_per_s: 0.0,
+        max_concurrent_requests,
+        steady_state: None,
+    }
+}
+
+/// Calculate multi-turn benchmark metrics with per-turn breakdown.
+///
+/// 1. Flatten all turns into SampleRequest/RequestFuncOutput pairs → overall metrics
+/// 2. Group by turn_index → per-turn metrics
+/// 3. Conversation-level stats
+pub fn calculate_multi_turn_metrics(
+    conversation_outputs: &[ConversationOutput],
+    dur_s: f64,
+    selected_percentiles: &[f64],
+    goodput_config: &GoodputConfig,
+    max_turn_count: usize,
+) -> MultiTurnMetrics {
+    // Flatten all turns for overall metrics
+    let mut all_requests: Vec<SampleRequest> = Vec::new();
+    let mut all_outputs: Vec<RequestFuncOutput> = Vec::new();
+
+    for conv in conversation_outputs {
+        for turn in &conv.turns {
+            all_requests.push(SampleRequest {
+                prompt: std::sync::Arc::from(""),
+                prompt_len: turn.cumulative_input_tokens,
+                expected_output_len: turn.request_output.output_tokens,
+                request_id: None,
+                ..Default::default()
+            });
+            all_outputs.push(turn.request_output.clone());
+        }
+    }
+
+    let (overall, _) = calculate_metrics(
+        &all_requests,
+        &all_outputs,
+        dur_s,
+        selected_percentiles,
+        true,
+        goodput_config,
+    );
+
+    // Per-turn breakdown
+    let mut per_turn: Vec<BenchmarkMetrics> = Vec::new();
+    for turn_idx in 0..max_turn_count {
+        let mut turn_requests: Vec<SampleRequest> = Vec::new();
+        let mut turn_outputs: Vec<RequestFuncOutput> = Vec::new();
+
+        for conv in conversation_outputs {
+            if let Some(turn) = conv.turns.get(turn_idx) {
+                turn_requests.push(SampleRequest {
+                    prompt: std::sync::Arc::from(""),
+                    prompt_len: turn.cumulative_input_tokens,
+                    expected_output_len: turn.request_output.output_tokens,
+                    request_id: None,
+                    ..Default::default()
+                });
+                turn_outputs.push(turn.request_output.clone());
+            }
+        }
+
+        if !turn_outputs.is_empty() {
+            let (turn_metrics, _) = calculate_metrics(
+                &turn_requests,
+                &turn_outputs,
+                dur_s,
+                selected_percentiles,
+                true,
+                goodput_config,
+            );
+            per_turn.push(turn_metrics);
+        }
+    }
+
+    // Conversation-level stats
+    let conversations_completed = conversation_outputs.iter().filter(|c| c.all_success).count();
+    let conversations_failed = conversation_outputs.len() - conversations_completed;
+
+    let avg_turns_completed = if conversation_outputs.is_empty() {
+        0.0
+    } else {
+        conversation_outputs.iter().map(|c| c.turns.len() as f64).sum::<f64>()
+            / conversation_outputs.len() as f64
+    };
+
+    let avg_conversation_duration_ms = if conversation_outputs.is_empty() {
+        0.0
+    } else {
+        conversation_outputs.iter().map(|c| c.total_duration_ms).sum::<f64>()
+            / conversation_outputs.len() as f64
+    };
+
+    MultiTurnMetrics {
+        overall,
+        per_turn,
+        conversations_completed,
+        conversations_failed,
+        avg_turns_completed,
+        avg_conversation_duration_ms,
+    }
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+
+    #[test]
+    fn test_mean_empty() {
+        assert_eq!(mean(&[]), 0.0);
+    }
+
+    #[test]
+    fn test_mean_values() {
+        assert!((mean(&[1.0, 2.0, 3.0]) - 2.0).abs() < 1e-10);
+    }
+
+    #[test]
+    fn test_median_odd() {
+        let sorted = sort_clone(&[3.0, 1.0, 2.0]);
+        assert!((median_sorted(&sorted) - 2.0).abs() < 1e-10);
+    }
+
+    #[test]
+    fn test_median_even() {
+        let sorted = sort_clone(&[1.0, 2.0, 3.0, 4.0]);
+        assert!((median_sorted(&sorted) - 2.5).abs() < 1e-10);
+    }
+
+    #[test]
+    fn test_percentile_99() {
+        let data: Vec<f64> = (0..100).map(|i| i as f64).collect();
+        let p99 = percentile_sorted(&data, 99.0);
+        assert!((p99 - 98.01).abs() < 0.1);
+    }
+}
@@ -0,0 +1,54 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+pub mod calculator;
+pub mod steady_state;
+
+use serde::{Deserialize, Serialize};
+pub use steady_state::SteadyStateMetrics;
+
+/// Multi-turn benchmark metrics with overall and per-turn breakdown.
+#[derive(Debug, Clone, Serialize, Deserialize)]
+pub struct MultiTurnMetrics {
+    pub overall: BenchmarkMetrics,
+    pub per_turn: Vec<BenchmarkMetrics>,
+    pub conversations_completed: usize,
+    pub conversations_failed: usize,
+    pub avg_turns_completed: f64,
+    pub avg_conversation_duration_ms: f64,
+}
+
+/// Full benchmark metrics for generation tasks.
+/// Matches Python's BenchmarkMetrics dataclass.
+#[derive(Debug, Clone, Serialize, Deserialize)]
+pub struct BenchmarkMetrics {
+    pub completed: usize,
+    pub failed: usize,
+    pub total_input: usize,
+    pub total_output: usize,
+    pub request_throughput: f64,
+    pub request_goodput: f64,
+    pub input_throughput: f64,
+    pub output_throughput: f64,
+    pub total_token_throughput: f64,
+    pub mean_ttft_ms: f64,
+    pub median_ttft_ms: f64,
+    pub std_ttft_ms: f64,
+    pub percentiles_ttft_ms: Vec<(f64, f64)>,
+    pub mean_tpot_ms: f64,
+    pub median_tpot_ms: f64,
+    pub std_tpot_ms: f64,
+    pub percentiles_tpot_ms: Vec<(f64, f64)>,
+    pub mean_itl_ms: f64,
+    pub median_itl_ms: f64,
+    pub std_itl_ms: f64,
+    pub percentiles_itl_ms: Vec<(f64, f64)>,
+    pub mean_e2el_ms: f64,
+    pub median_e2el_ms: f64,
+    pub std_e2el_ms: f64,
+    pub percentiles_e2el_ms: Vec<(f64, f64)>,
+    pub max_output_tokens_per_s: f64,
+    pub max_concurrent_requests: usize,
+    #[serde(default)]
+    pub steady_state: Option<SteadyStateMetrics>,
+}
@@ -0,0 +1,629 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+//! Steady-state metrics: throughput and TTFT measured over the window
+//! during which client-side in-flight concurrency is at or above a
+//! configurable fraction of `--max-concurrency`. Excludes ramp-up and
+//! drain phases to reduce run-to-run variance at very high concurrency.
+
+use serde::{Deserialize, Serialize};
+
+/// Bounds and metadata of the detected steady-state window.
+///
+/// `observed_peak` is event-exact and may differ by 1 from the
+/// bucket-approximated `max_concurrent_requests` on `BenchmarkMetrics`;
+/// this is expected.
+#[derive(Debug, Clone, Serialize, Deserialize)]
+pub struct SteadyStateWindow {
+    pub start_s: f64,
+    pub end_s: f64,
+    pub duration_s: f64,
+    pub target_concurrency: usize,
+    pub threshold: f64,
+    pub threshold_abs: usize,
+    pub observed_peak: usize,
+    pub requests_started_in_window: usize,
+    pub requests_completed_in_window: usize,
+    pub requests_total: usize,
+    pub warning: Option<String>,
+}
+
+#[derive(Debug, Clone, Serialize, Deserialize)]
+pub struct SteadyStateMetrics {
+    pub window: SteadyStateWindow,
+    pub request_throughput: f64,
+    pub input_throughput: f64,
+    pub output_throughput: f64,
+    pub total_token_throughput: f64,
+    pub mean_ttft_ms: f64,
+    pub median_ttft_ms: f64,
+    pub percentiles_ttft_ms: Vec<(f64, f64)>,
+    pub mean_tpot_ms: f64,
+    pub median_tpot_ms: f64,
+    pub p90_tpot_ms: f64,
+    pub p99_tpot_ms: f64,
+}
+
+use crate::backends::RequestFuncOutput;
+
+/// Detect the steady-state window over `outputs` under the given target concurrency.
+///
+/// Returns `None` when closed-loop gate is not met (no target), when concurrency
+/// never crosses `threshold_abs`, or when both started and completed in-window
+/// request counts are zero.
+///
+/// Timestamps in the returned window are normalized so `start_s` / `end_s` are
+/// "seconds from the earliest successful request's start_time."
+///
+/// `user_min_window_s`: if `Some`, sets the minimum window duration below which
+/// a `warning` is attached. If `None`, the caller is expected to have pre-resolved
+/// the default `max(10.0, 0.1 * total_run_duration_s)`.
+pub fn detect_window(
+    outputs: &[RequestFuncOutput],
+    target_concurrency: Option<usize>,
+    threshold: f64,
+    min_window_s: f64,
+    total_run_duration_s: f64,
+) -> Option<SteadyStateWindow> {
+    let target = target_concurrency?;
+    let threshold_abs = ((threshold * target as f64).ceil() as usize).max(1);
+
+    // Collect sorted start / end timestamps from successful requests only.
+    let mut starts: Vec<f64> = outputs.iter().filter(|o| o.success).map(|o| o.start_time).collect();
+    let mut ends: Vec<f64> =
+        outputs.iter().filter(|o| o.success).map(|o| o.start_time + o.latency).collect();
+    if starts.is_empty() {
+        return None;
+    }
+    starts.sort_by(|a, b| a.partial_cmp(b).unwrap_or(std::cmp::Ordering::Equal));
+    ends.sort_by(|a, b| a.partial_cmp(b).unwrap_or(std::cmp::Ordering::Equal));
+
+    let min_start = starts[0];
+    let requests_total = outputs.iter().filter(|o| o.success).count();
+
+    // Two-pointer merge. Tie rule: process `end` before `start` at equal timestamps.
+    let mut i = 0usize;
+    let mut j = 0usize;
+    let mut concurrency: usize = 0;
+    let mut observed_peak: usize = 0;
+    let mut up_crossing: Option<f64> = None;
+    let mut last_down_crossing: Option<f64> = None;
+
+    while i < starts.len() || j < ends.len() {
+        let next_is_end = match (starts.get(i), ends.get(j)) {
+            (Some(&s), Some(&e)) => e <= s, // tie -> end first
+            (None, Some(_)) => true,
+            (Some(_), None) => false,
+            (None, None) => break,
+        };
+
+        if next_is_end {
+            let t = ends[j];
+            j += 1;
+            let before = concurrency;
+            concurrency = concurrency.saturating_sub(1);
+            if before >= threshold_abs && concurrency < threshold_abs {
+                last_down_crossing = Some(t);
+            }
+        } else {
+            let t = starts[i];
+            i += 1;
+            let before = concurrency;
+            concurrency += 1;
+            if concurrency > observed_peak {
+                observed_peak = concurrency;
+            }
+            if before < threshold_abs && concurrency >= threshold_abs && up_crossing.is_none() {
+                up_crossing = Some(t);
+            }
+        }
+    }
+
+    let start_abs = up_crossing?;
+    // A down-crossing always exists once an up-crossing exists — the final ends
+    // drain concurrency to 0. Fall back defensively to the last end timestamp.
+    let end_abs = last_down_crossing.or_else(|| ends.last().copied())?;
+
+    let start_s = start_abs - min_start;
+    let end_s = end_abs - min_start;
+    let duration_s = end_s - start_s;
+
+    // Count requests whose start / completion events land in [start_abs, end_abs).
+    let mut requests_started_in_window = 0usize;
+    let mut requests_completed_in_window = 0usize;
+    for o in outputs.iter().filter(|o| o.success) {
+        if o.start_time >= start_abs && o.start_time < end_abs {
+            requests_started_in_window += 1;
+        }
+        let end_t = o.start_time + o.latency;
+        if end_t >= start_abs && end_t < end_abs {
+            requests_completed_in_window += 1;
+        }
+    }
+
+    if requests_started_in_window == 0 && requests_completed_in_window == 0 {
+        return None;
+    }
+
+    let warning = if duration_s < min_window_s {
+        Some(format!(
+            "steady-state window is {:.1}s ({:.1}% of run); may not be representative",
+            duration_s,
+            100.0 * duration_s / total_run_duration_s.max(1e-9)
+        ))
+    } else {
+        None
+    };
+
+    Some(SteadyStateWindow {
+        start_s,
+        end_s,
+        duration_s,
+        target_concurrency: target,
+        threshold,
+        threshold_abs,
+        observed_peak,
+        requests_started_in_window,
+        requests_completed_in_window,
+        requests_total,
+        warning,
+    })
+}
+
+use crate::datasets::SampleRequest;
+
+/// Compute steady-state metrics from per-request outputs, the detected window,
+/// and the original request list (for input-token attribution).
+///
+/// `is_pooling` gates off TTFT because pooling backends write `ttft = latency`
+/// as a placeholder; reporting TTFT over that would silently be E2EL.
+pub fn compute(
+    outputs: &[RequestFuncOutput],
+    input_requests: &[SampleRequest],
+    window: &SteadyStateWindow,
+    percentiles: &[f64],
+    is_pooling: bool,
+) -> SteadyStateMetrics {
+    // Re-derive min_start to convert normalized window bounds back to absolute times.
+    let min_start = outputs
+        .iter()
+        .filter(|o| o.success)
+        .map(|o| o.start_time)
+        .fold(f64::INFINITY, f64::min);
+    let start_abs = window.start_s + min_start;
+    let end_abs = window.end_s + min_start;
+    let duration = window.duration_s.max(1e-9);
+
+    let mut completed_in_window = 0usize;
+    let mut input_tokens_in_window: usize = 0;
+    let mut output_tokens_in_window: usize = 0;
+    let mut ttfts_in_window: Vec<f64> = Vec::new();
+    let mut tpots_in_window: Vec<f64> = Vec::new();
+
+    for (idx, o) in outputs.iter().enumerate() {
+        if !o.success {
+            continue;
+        }
+        let started_in = o.start_time >= start_abs && o.start_time < end_abs;
+        let end_t = o.start_time + o.latency;
+        let completed_in = end_t >= start_abs && end_t < end_abs;
+
+        if completed_in {
+            completed_in_window += 1;
+        }
+
+        if started_in {
+            input_tokens_in_window += input_requests[idx].prompt_len;
+            if !is_pooling {
+                ttfts_in_window.push(o.ttft);
+                // Per-request TPOT: (latency - ttft) / (output_tokens - 1).
+                // Matches calculator.rs:46-50; skip requests with <= 1 token.
+                if o.output_tokens > 1 {
+                    let tpot = (o.latency - o.ttft) / (o.output_tokens as f64 - 1.0);
+                    tpots_in_window.push(tpot);
+                }
+            }
+        }
+
+        // Per-token emission attribution (same approach as calculator.rs:98-113).
+        // Reconstruct: first token at start_time + ttft, then cumulative itl.
+        if o.output_tokens == 0 && o.itl.is_empty() {
+            continue;
+        }
+        let first_token_t = o.start_time + o.ttft;
+        if first_token_t >= start_abs && first_token_t < end_abs {
+            output_tokens_in_window += 1;
+        }
+        let mut t = first_token_t;
+        for &dt in &o.itl {
+            t += dt;
+            if t >= start_abs && t < end_abs {
+                output_tokens_in_window += 1;
+            }
+            if t >= end_abs {
+                break; // tokens only move forward in time
+            }
+        }
+    }
+
+    let request_throughput = completed_in_window as f64 / duration;
+    let input_throughput = input_tokens_in_window as f64 / duration;
+    let output_throughput = output_tokens_in_window as f64 / duration;
+    let total_token_throughput = input_throughput + output_throughput;
+
+    let (mean_ttft, median_ttft, pct_ttft) = if is_pooling {
+        (0.0, 0.0, Vec::new())
+    } else {
+        dist_stats(&ttfts_in_window, percentiles)
+    };
+
+    let (mean_tpot, median_tpot, pct_tpot) = if is_pooling {
+        (0.0, 0.0, Vec::new())
+    } else {
+        dist_stats(&tpots_in_window, &[90.0, 99.0])
+    };
+    let p90_tpot = pct_tpot.first().map(|(_, v)| *v).unwrap_or(0.0);
+    let p99_tpot = pct_tpot.get(1).map(|(_, v)| *v).unwrap_or(0.0);
+
+    SteadyStateMetrics {
+        window: window.clone(),
+        request_throughput,
+        input_throughput,
+        output_throughput,
+        total_token_throughput,
+        mean_ttft_ms: mean_ttft,
+        median_ttft_ms: median_ttft,
+        percentiles_ttft_ms: pct_ttft,
+        mean_tpot_ms: mean_tpot,
+        median_tpot_ms: median_tpot,
+        p90_tpot_ms: p90_tpot,
+        p99_tpot_ms: p99_tpot,
+    }
+}
+
+/// Returns (mean_ms, median_ms, percentiles_ms).  Values are seconds in; output
+/// is milliseconds.  Empty input yields zeros with a zero-filled percentile vec
+/// aligned to the requested percentile list.
+fn dist_stats(values: &[f64], percentiles: &[f64]) -> (f64, f64, Vec<(f64, f64)>) {
+    if values.is_empty() {
+        let pct = percentiles.iter().map(|&p| (p, 0.0)).collect();
+        return (0.0, 0.0, pct);
+    }
+    let mut sorted = values.to_vec();
+    sorted.sort_by(|a, b| a.partial_cmp(b).unwrap_or(std::cmp::Ordering::Equal));
+    let n = sorted.len();
+    let mean = sorted.iter().sum::<f64>() / n as f64;
+    let median = if n.is_multiple_of(2) {
+        (sorted[n / 2 - 1] + sorted[n / 2]) / 2.0
+    } else {
+        sorted[n / 2]
+    };
+    let pct = percentiles
+        .iter()
+        .map(|&p| {
+            let idx = (p / 100.0) * (n - 1) as f64;
+            let lo = idx.floor() as usize;
+            let hi = idx.ceil() as usize;
+            let frac = idx - lo as f64;
+            let v = if lo == hi {
+                sorted[lo]
+            } else {
+                sorted[lo] * (1.0 - frac) + sorted[hi] * frac
+            };
+            (p, v * 1000.0)
+        })
+        .collect();
+    (mean * 1000.0, median * 1000.0, pct)
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+    use crate::backends::RequestFuncOutput;
+
+    /// Build a successful synthetic output with given `start_time` and `latency`.
+    fn mk(start: f64, latency: f64) -> RequestFuncOutput {
+        RequestFuncOutput {
+            success: true,
+            start_time: start,
+            latency,
+            ttft: 0.0,
+            itl: Vec::new(),
+            output_tokens: 0,
+            ..Default::default()
+        }
+    }
+
+    #[test]
+    fn window_plateau_trapezoid() {
+        // 10 requests ramp in at t=0..10 each lasting 30s.
+        // Concurrency reaches 10 at t=10 and stays until t=30 when requests start ending.
+        // threshold = 1.0 means exact-target; threshold_abs = 10.
+        let mut outputs: Vec<RequestFuncOutput> = (0..10).map(|i| mk(i as f64, 30.0)).collect();
+        // Add one request fully inside the plateau so it can pick up as "in window"
+        outputs.push(mk(15.0, 5.0));
+
+        let w = detect_window(&outputs, Some(10), 1.0, 5.0, 30.0).unwrap();
+        assert_eq!(w.target_concurrency, 10);
+        assert_eq!(w.threshold_abs, 10);
+        // Up-crossing at t=9 (the 10th start) — concurrency transitions 9→10.
+        // Window start reported in normalized time (min_start = 0.0).
+        assert!((w.start_s - 9.0).abs() < 1e-9, "start_s = {}", w.start_s);
+        // Last down-crossing: the first end at t=30 (concurrency 10→9, falling below 10).
+        assert!((w.end_s - 30.0).abs() < 1e-9, "end_s = {}", w.end_s);
+        assert!(w.duration_s > 0.0);
+        assert_eq!(w.observed_peak, 11); // 10 ramp-in + 1 extra = 11 peak briefly
+        assert!(w.warning.is_none());
+    }
+
+    #[test]
+    fn window_none_when_target_none() {
+        let outputs = vec![mk(0.0, 1.0)];
+        assert!(detect_window(&outputs, None, 0.95, 1.0, 10.0).is_none());
+    }
+
+    #[test]
+    fn window_none_when_never_reaches_threshold() {
+        // Serial requests, concurrency never exceeds 1.
+        let outputs: Vec<RequestFuncOutput> = (0..5).map(|i| mk(i as f64 * 10.0, 1.0)).collect();
+        assert!(detect_window(&outputs, Some(5), 0.95, 1.0, 100.0).is_none());
+    }
+
+    #[test]
+    fn window_none_when_all_fail() {
+        let mut o = mk(0.0, 1.0);
+        o.success = false;
+        assert!(detect_window(&[o], Some(1), 1.0, 0.1, 10.0).is_none());
+    }
+
+    #[test]
+    fn window_threshold_one() {
+        // threshold=1.0, target=2 -> threshold_abs=2; needs both requests in flight.
+        let outputs = vec![mk(0.0, 5.0), mk(1.0, 5.0)];
+        let w = detect_window(&outputs, Some(2), 1.0, 0.1, 10.0).unwrap();
+        assert_eq!(w.threshold_abs, 2);
+        assert!((w.start_s - 1.0).abs() < 1e-9);
+        assert!((w.end_s - 5.0).abs() < 1e-9);
+    }
+
+    #[test]
+    fn window_ties_net_to_zero() {
+        // A request finishes and another starts at exactly the same instant;
+        // tie rule (end before start) prevents concurrency > 2 blip.
+        let outputs = vec![
+            mk(0.0, 5.0), // ends at 5.0
+            mk(2.0, 3.0), // ends at 5.0
+            mk(5.0, 2.0), // starts at 5.0 — ties with the two ends
+        ];
+        let w = detect_window(&outputs, Some(2), 1.0, 0.1, 20.0).unwrap();
+        assert_eq!(w.observed_peak, 2);
+    }
+
+    #[test]
+    fn window_warning_when_short() {
+        let outputs: Vec<RequestFuncOutput> = (0..10).map(|i| mk(i as f64 * 0.1, 1.0)).collect();
+        // Plateau is ~0.1s; min_window_s = 5s => warning.
+        let w = detect_window(&outputs, Some(10), 1.0, 5.0, 10.0).unwrap();
+        assert!(w.warning.is_some(), "expected warning");
+    }
+
+    #[test]
+    fn window_constant_concurrency_equals_full_run() {
+        // Concurrency is at target throughout — steady-state window covers the
+        // entire middle of the run where all 3 requests overlap.
+        let outputs = vec![mk(0.0, 10.0), mk(0.0, 10.0), mk(0.0, 10.0)];
+        let w = detect_window(&outputs, Some(3), 1.0, 0.1, 10.0).unwrap();
+        assert!((w.start_s - 0.0).abs() < 1e-9);
+        assert!((w.end_s - 10.0).abs() < 1e-9);
+    }
+
+    use std::sync::Arc;
+
+    use crate::datasets::SampleRequest;
+
+    fn mk_full(
+        start: f64,
+        latency: f64,
+        ttft: f64,
+        itl: Vec<f64>,
+        output_tokens: usize,
+    ) -> RequestFuncOutput {
+        RequestFuncOutput {
+            success: true,
+            start_time: start,
+            latency,
+            ttft,
+            itl,
+            output_tokens,
+            prompt_len: 10,
+            ..Default::default()
+        }
+    }
+
+    fn mk_sample(prompt_len: usize) -> SampleRequest {
+        SampleRequest {
+            prompt: Arc::from(""),
+            prompt_len,
+            expected_output_len: 0,
+            request_id: None,
+            ..Default::default()
+        }
+    }
+
+    #[test]
+    fn compute_constant_concurrency_matches_full_run() {
+        // 3 identical requests, each 10s, each emitting 5 tokens uniformly via itl=2s.
+        let itl = vec![2.0, 2.0, 2.0, 2.0]; // 5 tokens: first at t=start+ttft, then 4 more
+        let outputs = vec![
+            mk_full(0.0, 10.0, 2.0, itl.clone(), 5),
+            mk_full(0.0, 10.0, 2.0, itl.clone(), 5),
+            mk_full(0.0, 10.0, 2.0, itl.clone(), 5),
+        ];
+        let requests: Vec<SampleRequest> =
+            outputs.iter().map(|o| mk_sample(o.prompt_len)).collect();
+
+        let w = detect_window(&outputs, Some(3), 1.0, 0.1, 10.0).unwrap();
+        let m = compute(
+            &outputs,
+            &requests,
+            &w,
+            &[99.0],
+            false, // is_pooling
+        );
+
+        // Window is [0.0, 10.0). All 3 requests start in window (3).
+        // Tokens per request: first token at ttft=2s, then 4 more at itl=2s each
+        //   -> t = 2, 4, 6, 8, 10. The 10.0 token is AT end_s and is half-open excluded.
+        //   -> 4 tokens per request land in window.
+        assert_eq!(w.requests_started_in_window, 3);
+        // Half-open window [0.0, 10.0) excludes the end_t=10.0 completions, so
+        // request_throughput is exactly 0.0 here — tightens the half-open invariant.
+        assert!((m.request_throughput - 0.0).abs() < 1e-9);
+        // output tokens in window = 4 per request * 3 = 12, over 10s = 1.2 tok/s
+        assert!(
+            (m.output_throughput - 1.2).abs() < 1e-6,
+            "got {}",
+            m.output_throughput
+        );
+        // input throughput = 30 input tokens / 10s = 3.0
+        assert!(
+            (m.input_throughput - 3.0).abs() < 1e-9,
+            "got {}",
+            m.input_throughput
+        );
+    }
+
+    #[test]
+    fn compute_drain_only_tokens_not_counted() {
+        // 1 "steady" request defining the window, + 1 late-admitted request whose
+        // 1000 output tokens land after end_s — must not pump output_throughput.
+        let mut outputs = vec![mk_full(0.0, 10.0, 0.0, vec![1.0; 9], 10)];
+        // Late request: starts at 9.9 (within window), long output that drains past end_s.
+        let late_itl = vec![0.02; 999]; // 1000 tokens; emission starts after end_s
+        let late = mk_full(9.9, 30.0, 0.05, late_itl, 1000);
+        outputs.push(late);
+        let requests: Vec<SampleRequest> =
+            outputs.iter().map(|o| mk_sample(o.prompt_len)).collect();
+
+        // Target = 2 so threshold_abs = 2; window opens when both are in flight.
+        let w = detect_window(&outputs, Some(2), 1.0, 0.1, 40.0).unwrap();
+        let m = compute(&outputs, &requests, &w, &[99.0], false);
+
+        // Late request started in window but most of its tokens are emitted AFTER end_s.
+        // Only a handful (≤3) should count.
+        // Window end is 10.0 (the first request's completion).
+        // Late request first token at 9.9 + 0.05 = 9.95 (in window).
+        // Next tokens at 9.97, 9.99, 10.01... -> only ~3 fit before end_s=10.0.
+        let expected_max_late_tokens = 6.0; // generous bound
+        let window_dur = w.duration_s;
+        // Upper bound on output throughput: 10 (from req 1) + ≤6 (from late) / window_dur.
+        let upper_bound = (10.0 + expected_max_late_tokens) / window_dur;
+        assert!(
+            m.output_throughput <= upper_bound,
+            "output_throughput {} exceeds upper bound {} — drain tokens are being counted",
+            m.output_throughput,
+            upper_bound
+        );
+    }
+
+    #[test]
+    fn compute_pooling_skips_ttft_and_tpot() {
+        // For pooling backends, ttft == latency is a placeholder; we must NOT emit TTFT stats.
+        // Pooling has no output tokens so TPOT is also zeroed.
+        let outputs = vec![
+            mk_full(0.0, 5.0, 5.0, Vec::new(), 0),
+            mk_full(0.0, 5.0, 5.0, Vec::new(), 0),
+        ];
+        let requests: Vec<SampleRequest> =
+            outputs.iter().map(|o| mk_sample(o.prompt_len)).collect();
+        let w = detect_window(&outputs, Some(2), 1.0, 0.1, 10.0).unwrap();
+        let m = compute(&outputs, &requests, &w, &[99.0], /* is_pooling */ true);
+        assert_eq!(m.mean_ttft_ms, 0.0);
+        assert_eq!(m.median_ttft_ms, 0.0);
+        assert!(m.percentiles_ttft_ms.is_empty());
+        assert_eq!(m.mean_tpot_ms, 0.0);
+        assert_eq!(m.median_tpot_ms, 0.0);
+        assert_eq!(m.p90_tpot_ms, 0.0);
+        assert_eq!(m.p99_tpot_ms, 0.0);
+        // Throughput fields still populated.
+        assert!(m.input_throughput > 0.0);
+    }
+
+    #[test]
+    fn compute_tpot_matches_per_request_definition() {
+        // 3 requests, each 10s latency with ttft=2s, 5 output tokens.
+        // Per-request TPOT = (10 - 2) / (5 - 1) = 2s = 2000ms.
+        // All requests start at t=0 so all are in-window; window duration = 10s.
+        let itl = vec![2.0; 4];
+        let outputs = vec![
+            mk_full(0.0, 10.0, 2.0, itl.clone(), 5),
+            mk_full(0.0, 10.0, 2.0, itl.clone(), 5),
+            mk_full(0.0, 10.0, 2.0, itl.clone(), 5),
+        ];
+        let requests: Vec<SampleRequest> =
+            outputs.iter().map(|o| mk_sample(o.prompt_len)).collect();
+        let w = detect_window(&outputs, Some(3), 1.0, 0.1, 10.0).unwrap();
+        let m = compute(&outputs, &requests, &w, &[99.0], false);
+
+        assert!((m.mean_tpot_ms - 2000.0).abs() < 1e-6);
+        assert!((m.median_tpot_ms - 2000.0).abs() < 1e-6);
+        assert!((m.p90_tpot_ms - 2000.0).abs() < 1e-6);
+        assert!((m.p99_tpot_ms - 2000.0).abs() < 1e-6);
+    }
+
+    #[test]
+    fn json_roundtrip_some() {
+        let outputs = vec![mk_full(0.0, 10.0, 2.0, vec![2.0; 4], 5); 3];
+        let requests: Vec<SampleRequest> =
+            outputs.iter().map(|o| mk_sample(o.prompt_len)).collect();
+        let w = detect_window(&outputs, Some(3), 1.0, 0.1, 10.0).unwrap();
+        let m = compute(&outputs, &requests, &w, &[99.0], false);
+        let s = serde_json::to_string(&m).unwrap();
+        let parsed: SteadyStateMetrics = serde_json::from_str(&s).unwrap();
+        assert!((parsed.output_throughput - m.output_throughput).abs() < 1e-12);
+        assert_eq!(parsed.window.target_concurrency, 3);
+    }
+
+    #[test]
+    fn json_deserialize_missing_key_as_none() {
+        // Pre-feature JSON lacks `steady_state` — deserializing a wrapper that
+        // contains an Option field with #[serde(default)] should yield None.
+        #[derive(serde::Deserialize)]
+        struct Wrap {
+            #[serde(default)]
+            steady_state: Option<SteadyStateMetrics>,
+        }
+        let old_json = r#"{"other":"value"}"#;
+        let w: Wrap = serde_json::from_str(old_json).unwrap();
+        assert!(w.steady_state.is_none());
+    }
+
+    #[test]
+    fn end_to_end_trapezoid_run() {
+        // 5 requests ramp in at t=0..5, each lasts 20s, each emits 10 tokens
+        // at 1s intervals starting at ttft=1s.
+        let itl = vec![1.0; 9];
+        let outputs: Vec<RequestFuncOutput> =
+            (0..5).map(|i| mk_full(i as f64, 20.0, 1.0, itl.clone(), 10)).collect();
+        let requests: Vec<SampleRequest> = (0..5).map(|_| mk_sample(100)).collect();
+
+        let w = detect_window(&outputs, Some(5), 1.0, 1.0, 25.0).unwrap();
+        assert_eq!(w.threshold_abs, 5);
+        assert!((w.start_s - 4.0).abs() < 1e-9);
+        assert!((w.end_s - 20.0).abs() < 1e-9);
+
+        let m = compute(&outputs, &requests, &w, &[50.0, 99.0], false);
+
+        // Round-trip through JSON.
+        let s = serde_json::to_string(&m).unwrap();
+        let back: SteadyStateMetrics = serde_json::from_str(&s).unwrap();
+        assert!((back.request_throughput - m.request_throughput).abs() < 1e-12);
+        assert!((back.output_throughput - m.output_throughput).abs() < 1e-12);
+        assert_eq!(back.window.target_concurrency, 5);
+
+        // TTFT percentiles should have 2 entries matching the input percentiles.
+        assert_eq!(m.percentiles_ttft_ms.len(), 2);
+        assert!((m.percentiles_ttft_ms[0].0 - 50.0).abs() < 1e-9);
+        assert!((m.percentiles_ttft_ms[1].0 - 99.0).abs() < 1e-9);
+    }
+}
@@ -0,0 +1,357 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+use crate::config::BenchConfig;
+use crate::error::Result;
+
+/// Key metrics extracted from a single run's JSON result.
+struct RunMetrics {
+    request_throughput: f64,
+    output_throughput: f64,
+    total_token_throughput: f64,
+    mean_ttft_ms: f64,
+    median_ttft_ms: f64,
+    p99_ttft_ms: f64,
+    mean_tpot_ms: f64,
+    median_tpot_ms: f64,
+    p99_tpot_ms: f64,
+    mean_itl_ms: f64,
+    mean_e2el_ms: f64,
+    p99_e2el_ms: f64,
+    max_output_tokens_per_s: f64,
+    completed: f64,
+    failed: f64,
+    duration: f64,
+    ss_request_throughput: Option<f64>,
+    ss_output_throughput: Option<f64>,
+    ss_input_throughput: Option<f64>,
+    ss_total_token_throughput: Option<f64>,
+    ss_mean_ttft_ms: Option<f64>,
+    ss_median_ttft_ms: Option<f64>,
+    ss_mean_tpot_ms: Option<f64>,
+    ss_median_tpot_ms: Option<f64>,
+    ss_p90_tpot_ms: Option<f64>,
+    ss_p99_tpot_ms: Option<f64>,
+}
+
+impl RunMetrics {
+    fn from_json(json: &serde_json::Value) -> Self {
+        Self {
+            request_throughput: get(json, "request_throughput"),
+            output_throughput: get(json, "output_throughput"),
+            total_token_throughput: get(json, "total_token_throughput"),
+            mean_ttft_ms: get(json, "mean_ttft_ms"),
+            median_ttft_ms: get(json, "median_ttft_ms"),
+            p99_ttft_ms: get(json, "p99_ttft_ms"),
+            mean_tpot_ms: get(json, "mean_tpot_ms"),
+            median_tpot_ms: get(json, "median_tpot_ms"),
+            p99_tpot_ms: get(json, "p99_tpot_ms"),
+            mean_itl_ms: get(json, "mean_itl_ms"),
+            mean_e2el_ms: get(json, "mean_e2el_ms"),
+            p99_e2el_ms: get(json, "p99_e2el_ms"),
+            max_output_tokens_per_s: get(json, "max_output_tokens_per_s"),
+            completed: get(json, "completed"),
+            failed: get(json, "failed"),
+            duration: get(json, "duration"),
+            ss_request_throughput: get_ss_opt(json, "request_throughput"),
+            ss_output_throughput: get_ss_opt(json, "output_throughput"),
+            ss_input_throughput: get_ss_opt(json, "input_throughput"),
+            ss_total_token_throughput: get_ss_opt(json, "total_token_throughput"),
+            ss_mean_ttft_ms: get_ss_opt(json, "mean_ttft_ms"),
+            ss_median_ttft_ms: get_ss_opt(json, "median_ttft_ms"),
+            ss_mean_tpot_ms: get_ss_opt(json, "mean_tpot_ms"),
+            ss_median_tpot_ms: get_ss_opt(json, "median_tpot_ms"),
+            ss_p90_tpot_ms: get_ss_opt(json, "p90_tpot_ms"),
+            ss_p99_tpot_ms: get_ss_opt(json, "p99_tpot_ms"),
+        }
+    }
+}
+
+fn get(json: &serde_json::Value, key: &str) -> f64 {
+    json.get(key).and_then(|v| v.as_f64()).unwrap_or(0.0)
+}
+
+fn get_ss_opt(json: &serde_json::Value, key: &str) -> Option<f64> {
+    json.get("steady_state")
+        .and_then(|ss| if ss.is_null() { None } else { Some(ss) })
+        .and_then(|ss| ss.get(key))
+        .and_then(|v| v.as_f64())
+}
+
+/// Aggregated statistics for a single metric across N runs.
+struct MetricStats {
+    label: &'static str,
+    mean: f64,
+    std: f64,
+    min: f64,
+    max: f64,
+}
+
+/// Run the benchmark N times and report aggregated statistics.
+pub async fn run_multi(config: &BenchConfig, num_runs: usize) -> Result<()> {
+    println!(
+        "{:=^70}",
+        format!(" Multi-Run Benchmark ({num_runs} runs) ")
+    );
+    println!();
+
+    let mut all_runs: Vec<RunMetrics> = Vec::with_capacity(num_runs);
+
+    for i in 0..num_runs {
+        println!("{:-^70}", format!(" Run {}/{} ", i + 1, num_runs));
+
+        let mut run_config = config.clone();
+        // Suppress per-run save
+        run_config.save_result = false;
+        run_config.append_result = false;
+
+        let result = crate::benchmark::run_benchmark(&run_config).await?;
+        all_runs.push(RunMetrics::from_json(&result));
+
+        println!();
+    }
+
+    print_multi_run_summary(&all_runs);
+    Ok(())
+}
+
+fn print_multi_run_summary(runs: &[RunMetrics]) {
+    let n = runs.len();
+
+    // Collect each metric into a series, compute stats
+    let stats = vec![
+        compute_stats("Request throughput (req/s)", runs, |r| r.request_throughput),
+        compute_stats("Output throughput (tok/s)", runs, |r| r.output_throughput),
+        compute_stats("Total token throughput (tok/s)", runs, |r| {
+            r.total_token_throughput
+        }),
+        compute_stats("Peak output tokens/s", runs, |r| r.max_output_tokens_per_s),
+        compute_stats("Mean TTFT (ms)", runs, |r| r.mean_ttft_ms),
+        compute_stats("Median TTFT (ms)", runs, |r| r.median_ttft_ms),
+        compute_stats("P99 TTFT (ms)", runs, |r| r.p99_ttft_ms),
+        compute_stats("Mean TPOT (ms)", runs, |r| r.mean_tpot_ms),
+        compute_stats("Median TPOT (ms)", runs, |r| r.median_tpot_ms),
+        compute_stats("P99 TPOT (ms)", runs, |r| r.p99_tpot_ms),
+        compute_stats("Mean ITL (ms)", runs, |r| r.mean_itl_ms),
+        compute_stats("Mean E2EL (ms)", runs, |r| r.mean_e2el_ms),
+        compute_stats("P99 E2EL (ms)", runs, |r| r.p99_e2el_ms),
+        compute_stats("Completed requests", runs, |r| r.completed),
+        compute_stats("Failed requests", runs, |r| r.failed),
+        compute_stats("Duration (s)", runs, |r| r.duration),
+    ];
+
+    println!("{:=^80}", format!(" Multi-Run Summary ({n} runs) "));
+    println!(
+        "{:<35} {:>10} {:>10} {:>10} {:>10}",
+        "Metric", "Mean", "Std", "Min", "Max"
+    );
+    println!(
+        "{:-<35} {:->10} {:->10} {:->10} {:->10}",
+        "", "", "", "", ""
+    );
+
+    for s in &stats {
+        println!(
+            "{:<35} {:>10} {:>10} {:>10} {:>10}",
+            s.label,
+            fmt(s.mean),
+            fmt(s.std),
+            fmt(s.min),
+            fmt(s.max),
+        );
+    }
+
+    println!("{:=<80}", "");
+
+    // Print coefficient of variation for throughput
+    let tp = &stats[0]; // request_throughput
+    if tp.mean > 0.0 {
+        let cv = (tp.std / tp.mean) * 100.0;
+        println!(
+            "Throughput CV: {:.1}% (lower = more stable across runs)",
+            cv
+        );
+    }
+
+    // Steady-state summary — aggregate only runs that have a window.
+    let ss_stats = vec![
+        compute_stats_opt("Request throughput (req/s)", runs, |r| {
+            r.ss_request_throughput
+        }),
+        compute_stats_opt("Output throughput (tok/s)", runs, |r| {
+            r.ss_output_throughput
+        }),
+        compute_stats_opt("Input throughput (tok/s)", runs, |r| r.ss_input_throughput),
+        compute_stats_opt("Total token throughput (tok/s)", runs, |r| {
+            r.ss_total_token_throughput
+        }),
+        compute_stats_opt("Mean TTFT (ms)", runs, |r| r.ss_mean_ttft_ms),
+        compute_stats_opt("Median TTFT (ms)", runs, |r| r.ss_median_ttft_ms),
+        compute_stats_opt("Mean TPOT (ms)", runs, |r| r.ss_mean_tpot_ms),
+        compute_stats_opt("Median TPOT (ms)", runs, |r| r.ss_median_tpot_ms),
+        compute_stats_opt("P90 TPOT (ms)", runs, |r| r.ss_p90_tpot_ms),
+        compute_stats_opt("P99 TPOT (ms)", runs, |r| r.ss_p99_tpot_ms),
+    ];
+
+    let k_present = ss_stats[0].n_present;
+    let n_total = ss_stats[0].n_total;
+
+    if k_present == 0 {
+        println!();
+        println!("Steady-state: no runs had a valid window");
+    } else {
+        println!();
+        println!(
+            "{:=^80}",
+            format!(" Steady-State Summary ({k_present}/{n_total} runs) ")
+        );
+        if k_present < n_total {
+            println!(
+                "Aggregated over {k_present}/{n_total} runs ({} runs had no valid window)",
+                n_total - k_present
+            );
+        }
+        println!(
+            "{:<35} {:>10} {:>10} {:>10} {:>10}",
+            "Metric", "Mean", "Std", "Min", "Max"
+        );
+        for s in &ss_stats {
+            println!(
+                "{:<35} {:>10.2} {:>10.2} {:>10.2} {:>10.2}",
+                s.label, s.mean, s.std, s.min, s.max
+            );
+        }
+    }
+}
+
+fn compute_stats<F>(label: &'static str, runs: &[RunMetrics], f: F) -> MetricStats
+where
+    F: Fn(&RunMetrics) -> f64,
+{
+    let values: Vec<f64> = runs.iter().map(&f).collect();
+    let n = values.len() as f64;
+    let mean = values.iter().sum::<f64>() / n;
+    let variance = values.iter().map(|x| (x - mean).powi(2)).sum::<f64>() / n;
+    let std = variance.sqrt();
+    let min = values.iter().cloned().fold(f64::INFINITY, f64::min);
+    let max = values.iter().cloned().fold(f64::NEG_INFINITY, f64::max);
+
+    MetricStats {
+        label,
+        mean,
+        std,
+        min,
+        max,
+    }
+}
+
+fn fmt(v: f64) -> String {
+    if v == 0.0 {
+        "0".to_string()
+    } else if v == v.floor() && v.abs() < 1e9 {
+        format!("{}", v as i64)
+    } else {
+        format!("{:.2}", v)
+    }
+}
+
+/// Aggregated statistics over only the Some-valued runs.
+struct MetricStatsOpt {
+    label: &'static str,
+    mean: f64,
+    std: f64,
+    min: f64,
+    max: f64,
+    n_present: usize,
+    n_total: usize,
+}
+
+fn compute_stats_opt<F>(label: &'static str, runs: &[RunMetrics], f: F) -> MetricStatsOpt
+where
+    F: Fn(&RunMetrics) -> Option<f64>,
+{
+    let values: Vec<f64> = runs.iter().filter_map(&f).collect();
+    let n_present = values.len();
+    let n_total = runs.len();
+    if values.is_empty() {
+        return MetricStatsOpt {
+            label,
+            mean: 0.0,
+            std: 0.0,
+            min: 0.0,
+            max: 0.0,
+            n_present,
+            n_total,
+        };
+    }
+    let mean = values.iter().sum::<f64>() / n_present as f64;
+    let variance = values.iter().map(|x| (x - mean).powi(2)).sum::<f64>() / n_present as f64;
+    let std = variance.sqrt();
+    let min = values.iter().cloned().fold(f64::INFINITY, f64::min);
+    let max = values.iter().cloned().fold(f64::NEG_INFINITY, f64::max);
+    MetricStatsOpt {
+        label,
+        mean,
+        std,
+        min,
+        max,
+        n_present,
+        n_total,
+    }
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+
+    fn mk_run(ss: Option<f64>) -> RunMetrics {
+        RunMetrics {
+            request_throughput: 0.0,
+            output_throughput: 0.0,
+            total_token_throughput: 0.0,
+            mean_ttft_ms: 0.0,
+            median_ttft_ms: 0.0,
+            p99_ttft_ms: 0.0,
+            mean_tpot_ms: 0.0,
+            median_tpot_ms: 0.0,
+            p99_tpot_ms: 0.0,
+            mean_itl_ms: 0.0,
+            mean_e2el_ms: 0.0,
+            p99_e2el_ms: 0.0,
+            max_output_tokens_per_s: 0.0,
+            completed: 0.0,
+            failed: 0.0,
+            duration: 0.0,
+            ss_request_throughput: ss,
+            ss_output_throughput: None,
+            ss_input_throughput: None,
+            ss_total_token_throughput: None,
+            ss_mean_ttft_ms: None,
+            ss_median_ttft_ms: None,
+            ss_mean_tpot_ms: None,
+            ss_median_tpot_ms: None,
+            ss_p90_tpot_ms: None,
+            ss_p99_tpot_ms: None,
+        }
+    }
+
+    #[test]
+    fn compute_stats_opt_excludes_none() {
+        let runs = vec![mk_run(Some(10.0)), mk_run(None), mk_run(Some(20.0))];
+        let s = compute_stats_opt("x", &runs, |r| r.ss_request_throughput);
+        assert_eq!(s.n_present, 2);
+        assert_eq!(s.n_total, 3);
+        assert!((s.mean - 15.0).abs() < 1e-9);
+        assert_eq!(s.min, 10.0);
+        assert_eq!(s.max, 20.0);
+    }
+
+    #[test]
+    fn compute_stats_opt_all_none() {
+        let runs = vec![mk_run(None), mk_run(None)];
+        let s = compute_stats_opt("x", &runs, |r| r.ss_request_throughput);
+        assert_eq!(s.n_present, 0);
+        assert_eq!(s.n_total, 2);
+        assert_eq!(s.mean, 0.0);
+    }
+}
@@ -0,0 +1,827 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+use std::collections::HashMap;
+use std::sync::Arc;
+use std::time::Instant;
+
+use indicatif::{ProgressBar, ProgressStyle};
+use tokio::sync::Semaphore;
+
+use crate::backends::{Backend, RequestFuncInput, RequestFuncOutput, get_backend};
+use crate::benchmark::{
+    assign_lora_modules, compute_spec_decode_stats, fetch_spec_decode_metrics, pre_resolve_dns,
+    profile_on_batch_threshold, start_profiler_immediate, stop_profiler_immediate,
+};
+use crate::cli::DatasetName;
+use crate::config::BenchConfig;
+use crate::datasets::MultiTurnConversation;
+use crate::error::{BenchError, Result};
+use crate::metrics::calculator::calculate_multi_turn_metrics;
+use crate::output::console::print_multi_turn_results;
+use crate::output::json::{
+    append_result, build_multi_turn_result_json, compute_result_filename, save_result,
+};
+
+/// Output from a single turn within a conversation.
+#[derive(Debug, Clone)]
+pub struct TurnOutput {
+    pub turn_index: usize,
+    pub request_output: RequestFuncOutput,
+    pub cumulative_input_tokens: usize,
+}
+
+/// Output from an entire conversation.
+#[derive(Debug, Clone)]
+pub struct ConversationOutput {
+    pub conversation_id: String,
+    pub turns: Vec<TurnOutput>,
+    pub total_duration_ms: f64,
+    pub all_success: bool,
+}
+
+/// Run the multi-turn conversation benchmark.
+pub async fn run_multi_turn_benchmark(config: &BenchConfig) -> Result<serde_json::Value> {
+    let _ = get_backend(config.backend)?;
+
+    // Build HTTP client
+    let concurrency = config
+        .multi_turn_concurrency
+        .unwrap_or(config.max_concurrency.unwrap_or(config.num_prompts));
+
+    let mut client_builder = reqwest::Client::builder()
+        .pool_max_idle_per_host(concurrency.max(256))
+        .timeout(std::time::Duration::from_secs(6 * 60 * 60))
+        .connect_timeout(std::time::Duration::from_secs(30))
+        .tcp_keepalive(std::time::Duration::from_secs(60))
+        .tcp_nodelay(true)
+        .http1_only()
+        .no_proxy();
+
+    if config.insecure {
+        client_builder = client_builder.danger_accept_invalid_certs(true);
+    }
+
+    client_builder = pre_resolve_dns(&config.base_url, client_builder);
+
+    let client = client_builder
+        .build()
+        .map_err(|e| BenchError::Backend(format!("Failed to build HTTP client: {e}")))?;
+
+    // Resolve model
+    let (model_id, model_name) = if let Some(ref m) = config.model {
+        (m.clone(), config.model_name.clone())
+    } else {
+        println!("Model not specified, fetching first model from server...");
+        let (name, id) = get_first_model(&config.base_url, &client, &config.extra_headers).await?;
+        println!("First model name: {name}, first model id: {id}");
+        (id, Some(name))
+    };
+
+    // Load tokenizer
+    let tokenizer = if config.skip_tokenizer_init {
+        None
+    } else {
+        let tid = config.tokenizer_id.as_deref().unwrap_or(&model_id);
+        println!("Loading tokenizer: {tid}");
+        let server_info = Some((config.base_url.as_str(), model_id.as_str()));
+        let t = crate::tokenizer::load_tokenizer(tid, config.trust_remote_code, server_info)?;
+        println!("Tokenizer loaded successfully.");
+        Some(t)
+    };
+
+    // Generate/load conversations
+    println!("Generating multi-turn conversations...");
+    let gen_start = Instant::now();
+
+    let mut conversations = match config.dataset_name {
+        DatasetName::Random => {
+            let tok = tokenizer
+                .as_ref()
+                .ok_or_else(|| BenchError::Config("Random dataset requires a tokenizer".into()))?;
+            let prefix_sharing_config = if config.multi_turn_prefix_global_ratio > 0.0
+                || config.multi_turn_prefix_conversation_ratio > 0.0
+            {
+                Some(crate::datasets::multi_turn::PrefixSharingConfig {
+                    global_ratio: config.multi_turn_prefix_global_ratio,
+                    conversation_ratio: config.multi_turn_prefix_conversation_ratio,
+                })
+            } else {
+                None
+            };
+            let random_cfg = crate::datasets::multi_turn::MultiTurnRandomConfig {
+                num_conversations: config.num_prompts,
+                min_turns: config.multi_turn_min_turns,
+                max_turns: config.multi_turn_max_turns,
+                prefix_len: config.random_prefix_len,
+                input_len: config.random_input_len,
+                per_turn_input_len: config.per_turn_input_len,
+                output_len: config.random_output_len,
+                seed: config.seed,
+                request_id_prefix: config.request_id_prefix.clone(),
+                prefix_sharing_config,
+            };
+            crate::datasets::multi_turn::generate_multi_turn_random(tok, &random_cfg)?
+        }
+        DatasetName::ShareGpt => {
+            let tok = tokenizer.as_ref().ok_or_else(|| {
+                BenchError::Config("ShareGPT dataset requires a tokenizer".into())
+            })?;
+            let downloaded;
+            let path = match config.dataset_path.as_deref() {
+                Some(p) => p,
+                None => {
+                    downloaded = crate::datasets::sharegpt::download_sharegpt_dataset()?;
+                    downloaded.as_str()
+                }
+            };
+            crate::datasets::multi_turn::load_sharegpt_multi_turn(
+                tok,
+                path,
+                config.num_prompts,
+                config.sharegpt_output_len,
+                config.sharegpt_multi_turn_max_turns,
+                config.seed,
+                &config.request_id_prefix,
+            )?
+        }
+        DatasetName::RandomMm => {
+            return Err(BenchError::Config(
+                "Random-MM multi-turn is not yet supported. Use 'random' or 'sharegpt' with --multi-turn.".into(),
+            ));
+        }
+        DatasetName::Sonnet => {
+            return Err(BenchError::Config(
+                "Sonnet multi-turn is not yet supported. Use 'random' or 'sharegpt' with --multi-turn.".into(),
+            ));
+        }
+        DatasetName::SpeedBench => {
+            return Err(BenchError::Config(
+                "SPEED-Bench multi-turn is not yet supported. Use 'random' or 'sharegpt' with --multi-turn.".into(),
+            ));
+        }
+        DatasetName::Custom | DatasetName::PrefixRepetition | DatasetName::RandomRerank => {
+            return Err(BenchError::Config(
+                "This dataset does not support multi-turn. Use 'random' or 'sharegpt' with --multi-turn.".into(),
+            ));
+        }
+        DatasetName::Hf => {
+            return Err(BenchError::Config(
+                "HF dataset multi-turn is not yet supported. Use 'random' or 'sharegpt' with --multi-turn.".into(),
+            ));
+        }
+    };
+
+    let no_history = config.multi_turn_prefix_global_ratio > 0.0
+        || config.multi_turn_prefix_conversation_ratio > 0.0;
+
+    if let Some(max_model_len) = config.max_model_len {
+        let (filtered_conversations, filtered_turns) =
+            filter_turns_by_max_model_len(&mut conversations, max_model_len, no_history);
+        if filtered_turns > 0 || filtered_conversations > 0 {
+            println!(
+                "Filtered {filtered_turns} turn(s) and {filtered_conversations} conversation(s) above --max-model-len {max_model_len}."
+            );
+        }
+        if conversations.is_empty() {
+            return Err(BenchError::Config(
+                "No conversations remain after applying --max-model-len".into(),
+            ));
+        }
+    }
+
+    let gen_elapsed = gen_start.elapsed();
+    let total_turns: usize = conversations.iter().map(|c| c.turns.len()).sum();
+    println!(
+        "Generated {} conversations ({} total turns) in {:.2}s",
+        conversations.len(),
+        total_turns,
+        gen_elapsed.as_secs_f64()
+    );
+
+    // Log prefix sharing info
+    if no_history {
+        let num_special = tokenizer.as_ref().map(|t| t.num_special_tokens_to_add()).unwrap_or(0);
+        let real_input_len = config.random_input_len.saturating_sub(num_special);
+        let global_tokens =
+            (real_input_len as f64 * config.multi_turn_prefix_global_ratio).floor() as usize;
+        let conv_tokens =
+            (real_input_len as f64 * config.multi_turn_prefix_conversation_ratio).floor() as usize;
+        let unique_tokens = real_input_len.saturating_sub(global_tokens + conv_tokens);
+        println!(
+            "User message prefix sharing: {:.0}% global ({} tokens), {:.0}% per-conversation ({} tokens), {:.0}% unique ({} tokens)",
+            config.multi_turn_prefix_global_ratio * 100.0,
+            global_tokens,
+            config.multi_turn_prefix_conversation_ratio * 100.0,
+            conv_tokens,
+            (1.0 - config.multi_turn_prefix_global_ratio
+                - config.multi_turn_prefix_conversation_ratio)
+                * 100.0,
+            unique_tokens,
+        );
+        println!("No history accumulation: each turn sends fixed-length prompt only.");
+    }
+
+    if config.dry_run {
+        let total_user_tokens: usize = conversations
+            .iter()
+            .flat_map(|c| c.turns.iter())
+            .map(|t| t.user_message_len)
+            .sum();
+        println!("Dry run stats:");
+        println!("  Total conversations: {}", conversations.len());
+        println!("  Total turns: {total_turns}");
+        println!("  Total user message tokens: {total_user_tokens}");
+        return Ok(serde_json::json!({"dry_run": true, "mode": "multi_turn"}));
+    }
+
+    // Ready check with a simple single request
+    if config.ready_check_timeout_sec > 0 {
+        let first_turn = &conversations[0].turns[0];
+        let test_input = RequestFuncInput {
+            prompt: first_turn.user_message.clone(),
+            api_url: config.api_url.clone(),
+            prompt_len: first_turn.user_message_len,
+            output_len: first_turn.expected_output_len,
+            model: model_id.clone(),
+            model_name: model_name.clone(),
+            logprobs: config.logprobs,
+            extra_headers: config.extra_headers.clone(),
+            extra_body: config.extra_body.clone(),
+            ignore_eos: config.ignore_eos,
+            request_id: None,
+            ..Default::default()
+        };
+
+        println!("Starting initial single prompt test run...");
+        let test_output = crate::ready_checker::wait_for_endpoint(
+            config.backend,
+            &client,
+            &test_input,
+            config.ready_check_timeout_sec,
+            5,
+        )
+        .await?;
+        if !test_output.success {
+            return Err(BenchError::Backend(format!(
+                "Initial test run failed: {}",
+                test_output.error
+            )));
+        }
+        println!("Initial test run completed.");
+    }
+
+    // For random datasets in multi-turn mode, auto-set min_tokens to enforce
+    // output length without using ignore_eos (which causes unbounded context growth).
+    // min_tokens + max_completion_tokens together control output length precisely.
+    let extra_body = if config.dataset_name == DatasetName::Random && !config.ignore_eos {
+        let mut body = config.extra_body.clone().unwrap_or_else(|| serde_json::json!({}));
+        if let serde_json::Value::Object(ref mut map) = body
+            && !map.contains_key("min_tokens")
+        {
+            map.insert(
+                "min_tokens".to_string(),
+                serde_json::json!(config.random_output_len),
+            );
+            println!(
+                "Auto-setting min_tokens={} for multi-turn random dataset (use --extra-body to override)",
+                config.random_output_len
+            );
+        }
+        Some(body)
+    } else {
+        config.extra_body.clone()
+    };
+
+    // Fetch speculative decoding metrics before benchmark
+    let spec_decode_before =
+        fetch_spec_decode_metrics(&config.base_url, &client, &config.extra_headers).await;
+    if spec_decode_before.is_some() {
+        println!("Speculative decoding detected, will collect metrics.");
+    }
+
+    // Start profiler if requested (immediate mode — no batch threshold)
+    if config.profile && config.profile_batch_threshold.is_none() {
+        start_profiler_immediate(&client, &config.base_url, &config.extra_headers).await;
+    }
+
+    // Threshold-based profiling: spawn background task that polls /metrics
+    // and triggers start/stop profile when batch size is reached.
+    let profile_task = if let Some(threshold) = config.profile_batch_threshold {
+        let poll_client = client.clone();
+        let base_url = config.base_url.clone();
+        let extra_headers = config.extra_headers.clone();
+        let duration_secs = config.profile_duration;
+        let (cancel_tx, cancel_rx) = tokio::sync::oneshot::channel::<()>();
+        let handle = tokio::spawn(async move {
+            profile_on_batch_threshold(
+                &poll_client,
+                &base_url,
+                &extra_headers,
+                threshold,
+                duration_secs,
+                cancel_rx,
+            )
+            .await;
+        });
+        Some((cancel_tx, handle))
+    } else {
+        None
+    };
+
+    // Main benchmark
+    println!("Starting multi-turn benchmark...");
+    println!("Conversations: {}", conversations.len());
+    println!("Concurrency: {concurrency}");
+    println!("Inter-turn delay: {} ms", config.multi_turn_delay_ms);
+
+    let max_turn_count = conversations.iter().map(|c| c.turns.len()).max().unwrap_or(0);
+
+    // Progress bar counts total turns
+    let pb = if config.disable_tqdm {
+        None
+    } else {
+        let bar = ProgressBar::new(total_turns as u64);
+        bar.set_style(
+            ProgressStyle::with_template(
+                "{spinner:.green} [{elapsed_precise}] [{bar:40.cyan/blue}] {pos}/{len} turns ({eta})",
+            )
+            .unwrap()
+            .progress_chars("#>-"),
+        );
+        Some(bar)
+    };
+
+    let benchmark_start = Instant::now();
+
+    // Per-conversation LoRA assignment (sticky: all turns of a conversation
+    // share the same adapter). None when --lora-modules not set.
+    let lora_assignments = assign_lora_modules(
+        &config.lora_modules,
+        config.lora_assignment,
+        conversations.len(),
+        config.seed,
+    );
+    if let Some(modules) = config.lora_modules.as_ref() {
+        let names: Vec<&str> = modules.iter().map(|s| s.as_ref()).collect();
+        println!(
+            "LoRA adapters ({}): {:?} [assignment={:?}, scope=conversation]",
+            modules.len(),
+            names,
+            config.lora_assignment
+        );
+    }
+
+    // Shared state
+    let backend = get_backend(config.backend)?;
+    let api_url = Arc::new(config.api_url.clone());
+    let model = Arc::new(model_id.clone());
+    let model_name_arc = Arc::new(model_name.clone());
+    let extra_body = Arc::new(extra_body);
+    let base_extra_headers = Arc::new(config.extra_headers.clone());
+    let ignore_eos = config.ignore_eos;
+    let logprobs = config.logprobs;
+    let delay_ms = config.multi_turn_delay_ms;
+
+    // Semaphore controls max in-flight requests (not conversations).
+    // Each conversation acquires the permit only during the HTTP request,
+    // releasing it before the inter-turn delay so other conversations can
+    // immediately fill the slot.
+    let semaphore = Arc::new(Semaphore::new(concurrency));
+
+    // Spawn one task per conversation (all at once)
+    let mut handles = Vec::with_capacity(conversations.len());
+    for (i, conv) in conversations.into_iter().enumerate() {
+        let client = client.clone();
+        let backend = backend.clone();
+        let api_url = api_url.clone();
+        let model = model.clone();
+        let model_name = model_name_arc.clone();
+        let extra_body = extra_body.clone();
+        let base_extra_headers = base_extra_headers.clone();
+        let pb = pb.clone();
+        let bench_start = benchmark_start;
+        let semaphore = semaphore.clone();
+        let lora_name = lora_assignments.as_ref().map(|v| v[i].clone());
+
+        handles.push(tokio::spawn(async move {
+            run_conversation(
+                &conv,
+                &backend,
+                &client,
+                &api_url,
+                &model,
+                &model_name,
+                lora_name.as_deref(),
+                &extra_body,
+                &base_extra_headers,
+                ignore_eos,
+                logprobs,
+                delay_ms,
+                no_history,
+                bench_start,
+                pb.as_ref(),
+                &semaphore,
+            )
+            .await
+        }));
+    }
+
+    // Collect all conversation outputs
+    let mut all_outputs: Vec<ConversationOutput> = Vec::with_capacity(handles.len());
+    for handle in handles {
+        match handle.await {
+            Ok(output) => all_outputs.push(output),
+            Err(e) => {
+                eprintln!("Conversation task panicked: {e}");
+            }
+        }
+    }
+
+    if let Some(ref pb) = pb {
+        pb.finish_and_clear();
+    }
+
+    let benchmark_duration = benchmark_start.elapsed().as_secs_f64();
+
+    // Stop profiler if requested (immediate mode — no batch threshold)
+    if config.profile && config.profile_batch_threshold.is_none() {
+        stop_profiler_immediate(&client, &config.base_url, &config.extra_headers).await;
+    }
+
+    // Signal the threshold-based profile task that the benchmark is done, then wait
+    if let Some((cancel_tx, task)) = profile_task {
+        let _ = cancel_tx.send(());
+        if let Err(e) = task.await {
+            eprintln!("WARNING: Profile background task failed: {e}");
+        }
+    }
+
+    // Fetch speculative decoding metrics after benchmark and compute stats
+    let spec_decode_stats = if spec_decode_before.is_some() {
+        let spec_decode_after =
+            fetch_spec_decode_metrics(&config.base_url, &client, &config.extra_headers).await;
+        match (spec_decode_before.as_ref(), spec_decode_after.as_ref()) {
+            (Some(before), Some(after)) => compute_spec_decode_stats(before, after),
+            _ => None,
+        }
+    } else {
+        None
+    };
+
+    // Calculate metrics
+    let mt_metrics = calculate_multi_turn_metrics(
+        &all_outputs,
+        benchmark_duration,
+        &config.selected_percentiles,
+        &config.goodput,
+        max_turn_count,
+    );
+
+    // Print console output
+    print_multi_turn_results(
+        &mt_metrics,
+        benchmark_duration,
+        config,
+        spec_decode_stats.as_ref(),
+    );
+
+    // Build result JSON
+    let date_iso = chrono::Utc::now().format("%Y-%m-%dT%H:%M:%SZ").to_string();
+    let dt_filename = chrono::Local::now().format("%Y%m%d-%H%M%S").to_string();
+    let result_json = build_multi_turn_result_json(
+        config,
+        &mt_metrics,
+        &all_outputs,
+        benchmark_duration,
+        &date_iso,
+        spec_decode_stats.as_ref(),
+    );
+
+    // Save if requested
+    if config.save_result || config.append_result {
+        let model_for_filename = config.model.as_deref().unwrap_or(&model_id);
+        let file_name = compute_result_filename(config, model_for_filename, &dt_filename);
+
+        if let Some(ref dir) = config.result_dir {
+            std::fs::create_dir_all(dir)?;
+        }
+
+        if config.append_result {
+            append_result(&result_json, &file_name)?;
+        } else {
+            save_result(&result_json, &file_name)?;
+        }
+    }
+
+    Ok(result_json)
+}
+
+fn filter_turns_by_max_model_len(
+    conversations: &mut Vec<MultiTurnConversation>,
+    max_model_len: usize,
+    no_history: bool,
+) -> (usize, usize) {
+    let before_conversations = conversations.len();
+    let before_turns: usize = conversations.iter().map(|c| c.turns.len()).sum();
+
+    for conversation in conversations.iter_mut() {
+        if no_history {
+            conversation.turns.retain(|turn| {
+                turn.user_message_len.saturating_add(turn.expected_output_len) <= max_model_len
+            });
+            continue;
+        }
+
+        let keep_turns = valid_prefix_len_for_max_model_len(conversation, max_model_len);
+        conversation.turns.truncate(keep_turns);
+    }
+
+    conversations.retain(|conversation| conversation.turns.len() >= 2);
+
+    let after_conversations = conversations.len();
+    let after_turns: usize = conversations.iter().map(|c| c.turns.len()).sum();
+
+    (
+        before_conversations - after_conversations,
+        before_turns - after_turns,
+    )
+}
+
+fn valid_prefix_len_for_max_model_len(
+    conversation: &MultiTurnConversation,
+    max_model_len: usize,
+) -> usize {
+    let mut cumulative_input_tokens = 0usize;
+    let mut keep_turns = 0usize;
+
+    for turn in &conversation.turns {
+        cumulative_input_tokens = cumulative_input_tokens.saturating_add(turn.user_message_len);
+
+        if cumulative_input_tokens.saturating_add(turn.expected_output_len) > max_model_len {
+            break;
+        }
+
+        cumulative_input_tokens = cumulative_input_tokens.saturating_add(turn.expected_output_len);
+        keep_turns += 1;
+    }
+
+    keep_turns
+}
+
+/// Run a single conversation: sequential turns, building up message history.
+///
+/// `lora_name`, when set, replaces both `model` and `model_name` in every
+/// turn's request payload — sticky for the whole conversation so prefix-cache
+/// reuse across turns isn't broken by mid-conversation adapter switches.
+async fn run_conversation(
+    conversation: &MultiTurnConversation,
+    backend: &Backend,
+    client: &reqwest::Client,
+    api_url: &str,
+    model: &str,
+    model_name: &Option<String>,
+    lora_name: Option<&str>,
+    extra_body: &Option<serde_json::Value>,
+    base_extra_headers: &Option<HashMap<String, String>>,
+    ignore_eos: bool,
+    logprobs: Option<usize>,
+    delay_ms: u64,
+    no_history: bool,
+    bench_start: Instant,
+    pb: Option<&ProgressBar>,
+    semaphore: &Semaphore,
+) -> ConversationOutput {
+    let conv_start = Instant::now();
+    let mut messages: Vec<serde_json::Value> = Vec::new();
+    let mut cumulative_tokens: usize = 0;
+    let mut turn_outputs: Vec<TurnOutput> = Vec::new();
+    let mut all_success = true;
+
+    // Router affinity: all turns share same X-Session-ID
+    let mut extra_headers = base_extra_headers.clone().unwrap_or_default();
+    extra_headers.insert(
+        "X-Session-ID".to_string(),
+        conversation.conversation_id.clone(),
+    );
+
+    for (turn_idx, turn) in conversation.turns.iter().enumerate() {
+        if no_history {
+            // Prefix sharing mode: reset messages each turn, send only this turn's message
+            messages.clear();
+            messages.push(serde_json::json!({
+                "role": "user",
+                "content": [{"type": "text", "text": &*turn.user_message}]
+            }));
+            cumulative_tokens = turn.user_message_len;
+        } else {
+            // Normal mode: accumulate history
+            messages.push(serde_json::json!({
+                "role": "user",
+                "content": [{"type": "text", "text": &*turn.user_message}]
+            }));
+            cumulative_tokens += turn.user_message_len;
+        }
+
+        // Set min_tokens per-request so each turn's output length matches
+        // the dataset's expected length. Skip if ignore_eos is set (unbounded
+        // generation) or the user already provided min_tokens via --extra-body.
+        let turn_extra_body = if !ignore_eos {
+            let mut body = extra_body.clone().unwrap_or_else(|| serde_json::json!({}));
+            if let serde_json::Value::Object(ref mut map) = body
+                && !map.contains_key("min_tokens")
+            {
+                map.insert(
+                    "min_tokens".to_string(),
+                    serde_json::json!(turn.expected_output_len),
+                );
+            }
+            Some(body)
+        } else {
+            extra_body.clone()
+        };
+
+        let (req_model, req_model_name) = match lora_name {
+            Some(name) => (name.to_string(), Some(name.to_string())),
+            None => (model.to_string(), model_name.clone()),
+        };
+
+        let input = RequestFuncInput {
+            prompt: turn.user_message.clone(),
+            api_url: api_url.to_string(),
+            prompt_len: cumulative_tokens,
+            output_len: turn.expected_output_len,
+            model: req_model,
+            model_name: req_model_name,
+            logprobs,
+            extra_headers: Some(extra_headers.clone()),
+            extra_body: turn_extra_body,
+            ignore_eos,
+            request_id: Some(format!("{}-turn{}", conversation.conversation_id, turn_idx)),
+            messages: Some(serde_json::json!(messages)),
+            ..Default::default()
+        };
+
+        // Acquire semaphore permit before sending the request.
+        // Released after the response so the slot is free during inter-turn delay.
+        let _permit = match semaphore.acquire().await {
+            Ok(p) => p,
+            Err(_) => {
+                all_success = false;
+                break;
+            }
+        };
+        let request_instant = Instant::now();
+        let result = backend.send_request(&input, client).await;
+        drop(_permit);
+
+        let output = match result {
+            Ok(mut o) => {
+                o.start_time = request_instant.duration_since(bench_start).as_secs_f64();
+                o
+            }
+            Err(e) => RequestFuncOutput {
+                success: false,
+                error: e.to_string(),
+                prompt_len: cumulative_tokens,
+                start_time: request_instant.duration_since(bench_start).as_secs_f64(),
+                ..Default::default()
+            },
+        };
+
+        if let Some(pb) = pb {
+            pb.inc(1);
+        }
+
+        if output.success {
+            if !no_history {
+                // Add assistant response to history for next turn
+                messages.push(serde_json::json!({
+                    "role": "assistant",
+                    "content": &output.generated_text
+                }));
+                cumulative_tokens += output.output_tokens;
+            }
+
+            turn_outputs.push(TurnOutput {
+                turn_index: turn_idx,
+                request_output: output,
+                cumulative_input_tokens: cumulative_tokens,
+            });
+        } else {
+            all_success = false;
+            turn_outputs.push(TurnOutput {
+                turn_index: turn_idx,
+                request_output: output,
+                cumulative_input_tokens: cumulative_tokens,
+            });
+            break; // Conversation stops on failure
+        }
+
+        // Inter-turn delay (semaphore NOT held — slot is free for others)
+        if delay_ms > 0 && turn_idx + 1 < conversation.turns.len() {
+            tokio::time::sleep(std::time::Duration::from_millis(delay_ms)).await;
+        }
+    }
+
+    let total_duration_ms = conv_start.elapsed().as_secs_f64() * 1000.0;
+
+    ConversationOutput {
+        conversation_id: conversation.conversation_id.clone(),
+        turns: turn_outputs,
+        total_duration_ms,
+        all_success,
+    }
+}
+
+/// Fetch the first model from the server's /v1/models endpoint.
+async fn get_first_model(
+    base_url: &str,
+    client: &reqwest::Client,
+    extra_headers: &Option<HashMap<String, String>>,
+) -> Result<(String, String)> {
+    let url = format!("{base_url}/v1/models");
+    let mut request = client.get(&url);
+    if let Some(headers) = extra_headers {
+        for (k, v) in headers {
+            request = request.header(k, v);
+        }
+    }
+    if let Ok(api_key) = std::env::var("OPENAI_API_KEY") {
+        request = request.header("Authorization", format!("Bearer {api_key}"));
+    }
+
+    let response = request.send().await?;
+    let data: serde_json::Value = response.json().await?;
+
+    if let Some(models) = data.get("data").and_then(|d| d.as_array())
+        && let Some(first) = models.first()
+    {
+        let id = first.get("id").and_then(|v| v.as_str()).unwrap_or_default().to_string();
+        let root = first.get("root").and_then(|v| v.as_str()).unwrap_or(&id).to_string();
+        return Ok((id, root));
+    }
+
+    Err(BenchError::Config(format!(
+        "No models found on the server at {base_url}"
+    )))
+}
+
+#[cfg(test)]
+mod tests {
+    use std::sync::Arc;
+
+    use super::{filter_turns_by_max_model_len, valid_prefix_len_for_max_model_len};
+    use crate::datasets::{ConversationTurn, MultiTurnConversation};
+
+    fn conversation(turns: &[(usize, usize)]) -> MultiTurnConversation {
+        MultiTurnConversation {
+            conversation_id: "conv-0".to_string(),
+            turns: turns
+                .iter()
+                .map(|(user_message_len, expected_output_len)| ConversationTurn {
+                    user_message: Arc::from("hello"),
+                    user_message_len: *user_message_len,
+                    expected_output_len: *expected_output_len,
+                })
+                .collect(),
+        }
+    }
+
+    #[test]
+    fn test_valid_prefix_len_for_max_model_len_with_history() {
+        let conv = conversation(&[(40, 10), (45, 10)]);
+
+        assert_eq!(valid_prefix_len_for_max_model_len(&conv, 105), 2);
+        assert_eq!(valid_prefix_len_for_max_model_len(&conv, 104), 1);
+    }
+
+    #[test]
+    fn test_filter_turns_by_max_model_len_truncates_history_conversations() {
+        let mut conversations = vec![
+            conversation(&[(40, 10), (45, 10), (1, 1)]),
+            conversation(&[(100, 10), (1, 1)]),
+        ];
+
+        let (filtered_conversations, filtered_turns) =
+            filter_turns_by_max_model_len(&mut conversations, 105, false);
+
+        assert_eq!(filtered_conversations, 1);
+        assert_eq!(filtered_turns, 3);
+        assert_eq!(conversations.len(), 1);
+        assert_eq!(conversations[0].turns.len(), 2);
+    }
+
+    #[test]
+    fn test_filter_turns_by_max_model_len_retains_independent_no_history_turns() {
+        let mut conversations = vec![conversation(&[(40, 10), (95, 10), (1, 1)])];
+
+        let (filtered_conversations, filtered_turns) =
+            filter_turns_by_max_model_len(&mut conversations, 104, true);
+
+        assert_eq!(filtered_conversations, 0);
+        assert_eq!(filtered_turns, 1);
+        assert_eq!(conversations.len(), 1);
+        assert_eq!(conversations[0].turns.len(), 2);
+    }
+}
@@ -0,0 +1,400 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+use crate::benchmark::SpecDecodeStats;
+use crate::config::BenchConfig;
+use crate::metrics::{BenchmarkMetrics, MultiTurnMetrics};
+
+/// Print benchmark results to console in the same format as Python.
+///
+/// Mirrors the output format from serve.py:918-1093.
+pub fn print_results(
+    metrics: &BenchmarkMetrics,
+    benchmark_duration: f64,
+    config: &BenchConfig,
+    has_tokenizer: bool,
+    spec_decode_stats: Option<&SpecDecodeStats>,
+) {
+    let is_pooling = config.backend.is_pooling();
+
+    if is_pooling {
+        println!("{:=^60}", " Embedding/Pooling Benchmark Result ");
+    } else {
+        println!("{:=^60}", " Serving Benchmark Result ");
+    }
+    println!("{:<40} {:<10}", "Successful requests:", metrics.completed);
+    println!("{:<40} {:<10}", "Failed requests:", metrics.failed);
+
+    if let Some(mc) = config.max_concurrency {
+        println!("{:<40} {:<10}", "Maximum request concurrency:", mc);
+    }
+    if !config.request_rate.is_infinite() {
+        println!(
+            "{:<40} {:<10.2}",
+            "Request rate configured (RPS):", config.request_rate
+        );
+    }
+    println!(
+        "{:<40} {:<10.2}",
+        "Benchmark duration (s):", benchmark_duration
+    );
+    println!("{:<40} {:<10}", "Total input tokens:", metrics.total_input);
+
+    if has_tokenizer && !is_pooling {
+        println!(
+            "{:<40} {:<10}",
+            "Total generated tokens:", metrics.total_output
+        );
+    }
+
+    println!(
+        "{:<40} {:<10.2}",
+        "Request throughput (req/s):", metrics.request_throughput
+    );
+    if metrics.request_goodput > 0.0 {
+        println!(
+            "{:<40} {:<10.2}",
+            "Request goodput (req/s):", metrics.request_goodput
+        );
+    }
+
+    if !is_pooling {
+        if has_tokenizer {
+            println!(
+                "{:<40} {:<10.2}",
+                "Output token throughput (tok/s):", metrics.output_throughput
+            );
+            println!(
+                "{:<40} {:<10.2}",
+                "Peak output token throughput (tok/s):", metrics.max_output_tokens_per_s
+            );
+        }
+
+        println!(
+            "{:<40} {:<10.2}",
+            "Peak concurrent requests:", metrics.max_concurrent_requests as f64
+        );
+
+        if has_tokenizer {
+            println!(
+                "{:<40} {:<10.2}",
+                "Total token throughput (tok/s):", metrics.total_token_throughput
+            );
+        }
+    } else {
+        println!(
+            "{:<40} {:<10.2}",
+            "Input token throughput (tok/s):", metrics.input_throughput
+        );
+        println!(
+            "{:<40} {:<10.2}",
+            "Peak concurrent requests:", metrics.max_concurrent_requests as f64
+        );
+    }
+
+    // Print per-metric percentiles
+    if has_tokenizer && !is_pooling {
+        print_metric_section(
+            "ttft",
+            "TTFT",
+            "Time to First Token",
+            &config.selected_percentile_metrics,
+            metrics.mean_ttft_ms,
+            metrics.median_ttft_ms,
+            &metrics.percentiles_ttft_ms,
+        );
+        print_metric_section(
+            "tpot",
+            "TPOT",
+            "Time per Output Token (excl. 1st token)",
+            &config.selected_percentile_metrics,
+            metrics.mean_tpot_ms,
+            metrics.median_tpot_ms,
+            &metrics.percentiles_tpot_ms,
+        );
+        print_metric_section(
+            "itl",
+            "ITL",
+            "Inter-token Latency",
+            &config.selected_percentile_metrics,
+            metrics.mean_itl_ms,
+            metrics.median_itl_ms,
+            &metrics.percentiles_itl_ms,
+        );
+    }
+    print_metric_section(
+        "e2el",
+        "E2EL",
+        "End-to-end Latency",
+        &config.selected_percentile_metrics,
+        metrics.mean_e2el_ms,
+        metrics.median_e2el_ms,
+        &metrics.percentiles_e2el_ms,
+    );
+
+    if let Some(stats) = spec_decode_stats {
+        print_spec_decode_section(stats);
+    }
+
+    println!("{:=<60}", "");
+
+    print_steady_state(metrics);
+}
+
+/// Print multi-turn benchmark results to console.
+pub fn print_multi_turn_results(
+    mt_metrics: &MultiTurnMetrics,
+    benchmark_duration: f64,
+    config: &BenchConfig,
+    spec_decode_stats: Option<&SpecDecodeStats>,
+) {
+    println!("{:=^60}", " Multi-Turn Benchmark Result ");
+    println!(
+        "{:<45} {}/{}",
+        "Conversations completed/total:",
+        mt_metrics.conversations_completed,
+        mt_metrics.conversations_completed + mt_metrics.conversations_failed,
+    );
+    println!(
+        "{:<45} {}",
+        "Turns per conversation (configured):", config.multi_turn_num_turns,
+    );
+    println!(
+        "{:<45} {:.1}",
+        "Avg turns completed:", mt_metrics.avg_turns_completed,
+    );
+    println!(
+        "{:<45} {:.0}",
+        "Avg conversation duration (ms):", mt_metrics.avg_conversation_duration_ms,
+    );
+    println!(
+        "{:<45} {}",
+        "Concurrency:",
+        config
+            .multi_turn_concurrency
+            .unwrap_or(config.max_concurrency.unwrap_or(config.num_prompts)),
+    );
+    println!(
+        "{:<45} {}",
+        "Inter-turn delay (ms):", config.multi_turn_delay_ms,
+    );
+    println!(
+        "{:<45} {:.2}",
+        "Benchmark duration (s):", benchmark_duration,
+    );
+
+    // Overall metrics
+    println!("{:-^60}", " Overall (All Turns) ");
+    print_metrics_block(&mt_metrics.overall, config);
+
+    // Per-turn breakdown
+    for (i, turn_metrics) in mt_metrics.per_turn.iter().enumerate() {
+        let samples = turn_metrics.completed + turn_metrics.failed;
+        println!("{:-^60}", format!(" Turn {} ({} samples) ", i + 1, samples));
+        print_metrics_block(turn_metrics, config);
+    }
+
+    if let Some(stats) = spec_decode_stats {
+        print_spec_decode_section(stats);
+    }
+
+    println!("{:=<60}", "");
+}
+
+/// Print a metrics block (reused for overall and per-turn).
+fn print_metrics_block(metrics: &BenchmarkMetrics, config: &BenchConfig) {
+    println!("{:<45} {:<10}", "Successful requests:", metrics.completed);
+    println!("{:<45} {:<10}", "Failed requests:", metrics.failed);
+    println!("{:<45} {:<10}", "Total input tokens:", metrics.total_input);
+    println!(
+        "{:<45} {:<10}",
+        "Total generated tokens:", metrics.total_output
+    );
+    println!(
+        "{:<45} {:<10.2}",
+        "Request throughput (req/s):", metrics.request_throughput
+    );
+    println!(
+        "{:<45} {:<10.2}",
+        "Input token throughput (tok/s):", metrics.input_throughput
+    );
+    println!(
+        "{:<45} {:<10.2}",
+        "Output token throughput (tok/s):", metrics.output_throughput
+    );
+    println!(
+        "{:<45} {:<10.2}",
+        "Total token throughput (tok/s):", metrics.total_token_throughput
+    );
+
+    print_metric_section(
+        "ttft",
+        "TTFT",
+        "Time to First Token",
+        &config.selected_percentile_metrics,
+        metrics.mean_ttft_ms,
+        metrics.median_ttft_ms,
+        &metrics.percentiles_ttft_ms,
+    );
+    print_metric_section(
+        "tpot",
+        "TPOT",
+        "Time per Output Token (excl. 1st token)",
+        &config.selected_percentile_metrics,
+        metrics.mean_tpot_ms,
+        metrics.median_tpot_ms,
+        &metrics.percentiles_tpot_ms,
+    );
+    print_metric_section(
+        "itl",
+        "ITL",
+        "Inter-token Latency",
+        &config.selected_percentile_metrics,
+        metrics.mean_itl_ms,
+        metrics.median_itl_ms,
+        &metrics.percentiles_itl_ms,
+    );
+    print_metric_section(
+        "e2el",
+        "E2EL",
+        "End-to-end Latency",
+        &config.selected_percentile_metrics,
+        metrics.mean_e2el_ms,
+        metrics.median_e2el_ms,
+        &metrics.percentiles_e2el_ms,
+    );
+}
+
+fn print_metric_section(
+    attr_name: &str,
+    short_name: &str,
+    header: &str,
+    selected: &[String],
+    mean_ms: f64,
+    median_ms: f64,
+    percentiles: &[(f64, f64)],
+) {
+    if !selected.iter().any(|s| s == attr_name) {
+        return;
+    }
+    println!("{:-^60}", header);
+    println!(
+        "{:<40} {:<10.2}",
+        format!("Mean {short_name} (ms):"),
+        mean_ms
+    );
+    println!(
+        "{:<40} {:<10.2}",
+        format!("Median {short_name} (ms):"),
+        median_ms
+    );
+    for (p, value) in percentiles {
+        let p_str = if *p == p.floor() {
+            format!("{}", *p as i64)
+        } else {
+            format!("{p}")
+        };
+        println!(
+            "{:<40} {:<10.2}",
+            format!("P{p_str} {short_name} (ms):"),
+            value
+        );
+    }
+}
+
+/// Print speculative decoding metrics section.
+fn print_spec_decode_section(stats: &SpecDecodeStats) {
+    println!("{:-^50}", "Speculative Decoding");
+    println!(
+        "{:<40} {:<10.2}",
+        "Acceptance rate (%):", stats.acceptance_rate
+    );
+    println!(
+        "{:<40} {:<10.2}",
+        "Acceptance length:", stats.acceptance_length
+    );
+    println!("{:<40} {:<10}", "Drafts:", stats.num_drafts);
+    println!("{:<40} {:<10}", "Draft tokens:", stats.draft_tokens);
+    println!("{:<40} {:<10}", "Accepted tokens:", stats.accepted_tokens);
+    if !stats.per_position_acceptance_rates.is_empty() {
+        println!("Per-position acceptance (%):");
+        for (i, rate) in stats.per_position_acceptance_rates.iter().enumerate() {
+            println!("{:<40} {:<10.2}", format!("  Position {i}:"), rate * 100.0);
+        }
+    }
+}
+
+fn print_steady_state(metrics: &crate::metrics::BenchmarkMetrics) {
+    let Some(ss) = metrics.steady_state.as_ref() else {
+        return;
+    };
+
+    if let Some(warning) = ss.window.warning.as_ref() {
+        println!("Warning: {warning}");
+    }
+
+    let total = ss.window.requests_total.max(1);
+    let started_pct = 100.0 * ss.window.requests_started_in_window as f64 / total as f64;
+
+    println!("{:=^47}", " Steady-State Metrics ");
+    println!(
+        "{:<35} >= {:.2} * {} = {}",
+        "Concurrency threshold:",
+        ss.window.threshold,
+        ss.window.target_concurrency,
+        ss.window.threshold_abs,
+    );
+    println!(
+        "{:<35} {:.1}s -> {:.1}s  ({:.1}s)",
+        "Window:", ss.window.start_s, ss.window.end_s, ss.window.duration_s
+    );
+    println!(
+        "{:<35} {}",
+        "Observed peak concurrency:", ss.window.observed_peak
+    );
+    println!(
+        "{:<35} {} / {} ({:.1}%)",
+        "Requests started in window:",
+        ss.window.requests_started_in_window,
+        ss.window.requests_total,
+        started_pct
+    );
+    println!(
+        "{:<35} {}",
+        "Requests completed in window:", ss.window.requests_completed_in_window
+    );
+    println!("{:-<47}", "");
+    println!(
+        "{:<35} {:.2}",
+        "Request throughput (req/s):", ss.request_throughput
+    );
+    println!(
+        "{:<35} {:.2}",
+        "Output token throughput (tok/s):", ss.output_throughput
+    );
+    println!(
+        "{:<35} {:.2}",
+        "Input token throughput (tok/s):", ss.input_throughput
+    );
+    println!(
+        "{:<35} {:.2}",
+        "Total token throughput (tok/s):", ss.total_token_throughput
+    );
+
+    if !ss.percentiles_ttft_ms.is_empty() || ss.mean_ttft_ms > 0.0 {
+        println!("{:-<47}", "");
+        println!("{:<35} {:.2}", "Mean TTFT (ms):", ss.mean_ttft_ms);
+        println!("{:<35} {:.2}", "Median TTFT (ms):", ss.median_ttft_ms);
+        for (p, v) in &ss.percentiles_ttft_ms {
+            println!("{:<35} {:.2}", format!("P{} TTFT (ms):", *p as u32), v);
+        }
+    }
+    if ss.mean_tpot_ms > 0.0 || ss.median_tpot_ms > 0.0 {
+        println!("{:-<47}", "");
+        println!("{:<35} {:.2}", "Mean TPOT (ms):", ss.mean_tpot_ms);
+        println!("{:<35} {:.2}", "Median TPOT (ms):", ss.median_tpot_ms);
+        println!("{:<35} {:.2}", "P90 TPOT (ms):", ss.p90_tpot_ms);
+        println!("{:<35} {:.2}", "P99 TPOT (ms):", ss.p99_tpot_ms);
+    }
+    println!("{:=^47}", "");
+}
@@ -0,0 +1,776 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+use serde_json::Value;
+
+use crate::backends::RequestFuncOutput;
+use crate::benchmark::SpecDecodeStats;
+use crate::config::BenchConfig;
+use crate::error::Result;
+use crate::metrics::{BenchmarkMetrics, MultiTurnMetrics};
+use crate::multi_turn::ConversationOutput;
+
+/// Build the result JSON object matching the Python output schema exactly.
+///
+/// Mirrors serve.py:1801-1947 result_json construction.
+pub fn build_result_json(
+    config: &BenchConfig,
+    metrics: &BenchmarkMetrics,
+    actual_output_lens: &[usize],
+    outputs: &[RequestFuncOutput],
+    benchmark_duration: f64,
+    current_dt: &str,
+    spec_decode_stats: Option<&SpecDecodeStats>,
+) -> serde_json::Value {
+    let mut result = serde_json::Map::new();
+
+    // Setup
+    result.insert("date".into(), Value::String(current_dt.to_string()));
+    result.insert(
+        "endpoint_type".into(),
+        Value::String(config.backend.to_string()),
+    );
+    result.insert("backend".into(), Value::String(config.backend.to_string()));
+    result.insert(
+        "label".into(),
+        config.label.as_ref().map(|l| Value::String(l.clone())).unwrap_or(Value::Null),
+    );
+    result.insert(
+        "model_id".into(),
+        Value::String(config.model.clone().unwrap_or_default()),
+    );
+    result.insert(
+        "tokenizer_id".into(),
+        config
+            .tokenizer_id
+            .as_ref()
+            .map(|t| Value::String(t.clone()))
+            .unwrap_or(Value::Null),
+    );
+    result.insert(
+        "num_prompts".into(),
+        Value::Number(config.num_prompts.into()),
+    );
+    result.insert(
+        "max_model_len".into(),
+        config.max_model_len.map(|v| Value::Number(v.into())).unwrap_or(Value::Null),
+    );
+
+    // Metadata
+    if let Some(ref metadata) = config.metadata {
+        for (k, v) in metadata {
+            result.insert(k.clone(), Value::String(v.clone()));
+        }
+    }
+
+    // Traffic
+    if config.request_rate.is_infinite() {
+        result.insert("request_rate".into(), Value::String("inf".to_string()));
+    } else {
+        result.insert(
+            "request_rate".into(),
+            serde_json::json!(config.request_rate),
+        );
+    }
+    result.insert("burstiness".into(), serde_json::json!(config.burstiness));
+    result.insert(
+        "max_concurrency".into(),
+        config.max_concurrency.map(|v| Value::Number(v.into())).unwrap_or(Value::Null),
+    );
+
+    let is_pooling = config.backend.is_pooling();
+
+    // Benchmark results
+    result.insert("duration".into(), serde_json::json!(benchmark_duration));
+    result.insert("completed".into(), serde_json::json!(metrics.completed));
+    result.insert("failed".into(), serde_json::json!(metrics.failed));
+    result.insert(
+        "total_input_tokens".into(),
+        serde_json::json!(metrics.total_input),
+    );
+    if !is_pooling {
+        result.insert(
+            "total_output_tokens".into(),
+            serde_json::json!(metrics.total_output),
+        );
+    }
+    result.insert(
+        "request_throughput".into(),
+        serde_json::json!(metrics.request_throughput),
+    );
+    if config.goodput.is_empty() {
+        result.insert("request_goodput".into(), Value::Null);
+    } else {
+        result.insert(
+            "request_goodput".into(),
+            serde_json::json!(metrics.request_goodput),
+        );
+    }
+    if is_pooling {
+        result.insert(
+            "total_token_throughput".into(),
+            serde_json::json!(metrics.total_token_throughput),
+        );
+    } else {
+        result.insert(
+            "input_throughput".into(),
+            serde_json::json!(metrics.input_throughput),
+        );
+        result.insert(
+            "output_throughput".into(),
+            serde_json::json!(metrics.output_throughput),
+        );
+        result.insert(
+            "total_token_throughput".into(),
+            serde_json::json!(metrics.total_token_throughput),
+        );
+        result.insert(
+            "max_output_tokens_per_s".into(),
+            serde_json::json!(metrics.max_output_tokens_per_s),
+        );
+        result.insert(
+            "max_concurrent_requests".into(),
+            serde_json::json!(metrics.max_concurrent_requests),
+        );
+        // Inverse Real-Time Factor (for ASR benchmarks; 0.0 for generation)
+        result.insert("rtfx".into(), serde_json::json!(0.0));
+    }
+
+    // Per-failure log (always emitted when there are failures, regardless of --save-detailed).
+    // ttft/itl/output_tokens/start_time/latency are best-effort — zero/empty for failures
+    // that occurred before any tokens were received. Non-zero values indicate a partial
+    // failure (some tokens streamed before the error).
+    if metrics.failed > 0 {
+        result.insert(
+            "failed_requests".into(),
+            Value::Array(collect_failed_requests(outputs)),
+        );
+    }
+
+    // Per-request data.
+    // For pooling: Python always includes input_lens and errors (serve.py:1004-1013).
+    // For generation: all per-request arrays are gated by --save-detailed.
+    if is_pooling {
+        result.insert(
+            "input_lens".into(),
+            serde_json::json!(outputs.iter().map(|o| o.prompt_len).collect::<Vec<_>>()),
+        );
+        result.insert(
+            "errors".into(),
+            serde_json::json!(outputs.iter().map(|o| &o.error).collect::<Vec<_>>()),
+        );
+    }
+    if config.save_detailed {
+        if !is_pooling {
+            result.insert(
+                "input_lens".into(),
+                serde_json::json!(outputs.iter().map(|o| o.prompt_len).collect::<Vec<_>>()),
+            );
+            result.insert("output_lens".into(), serde_json::json!(actual_output_lens));
+            result.insert(
+                "ttfts".into(),
+                serde_json::json!(outputs.iter().map(|o| o.ttft).collect::<Vec<_>>()),
+            );
+            result.insert(
+                "itls".into(),
+                serde_json::json!(outputs.iter().map(|o| &o.itl).collect::<Vec<_>>()),
+            );
+            result.insert(
+                "generated_texts".into(),
+                serde_json::json!(outputs.iter().map(|o| &o.generated_text).collect::<Vec<_>>()),
+            );
+            result.insert(
+                "errors".into(),
+                serde_json::json!(outputs.iter().map(|o| &o.error).collect::<Vec<_>>()),
+            );
+        }
+        result.insert(
+            "latencies".into(),
+            serde_json::json!(outputs.iter().map(|o| o.latency).collect::<Vec<_>>()),
+        );
+        result.insert(
+            "start_times".into(),
+            serde_json::json!(outputs.iter().map(|o| o.start_time).collect::<Vec<_>>()),
+        );
+    }
+
+    // Speculative decoding stats
+    if let Some(stats) = spec_decode_stats {
+        insert_spec_decode_stats(&mut result, stats);
+    }
+
+    // Per-metric stats (pooling only has e2el)
+    if !is_pooling {
+        add_metric_stats(
+            &mut result,
+            "ttft",
+            metrics.mean_ttft_ms,
+            metrics.median_ttft_ms,
+            metrics.std_ttft_ms,
+            &metrics.percentiles_ttft_ms,
+            &config.selected_percentile_metrics,
+        );
+        add_metric_stats(
+            &mut result,
+            "tpot",
+            metrics.mean_tpot_ms,
+            metrics.median_tpot_ms,
+            metrics.std_tpot_ms,
+            &metrics.percentiles_tpot_ms,
+            &config.selected_percentile_metrics,
+        );
+        add_metric_stats(
+            &mut result,
+            "itl",
+            metrics.mean_itl_ms,
+            metrics.median_itl_ms,
+            metrics.std_itl_ms,
+            &metrics.percentiles_itl_ms,
+            &config.selected_percentile_metrics,
+        );
+    }
+    add_metric_stats(
+        &mut result,
+        "e2el",
+        metrics.mean_e2el_ms,
+        metrics.median_e2el_ms,
+        metrics.std_e2el_ms,
+        &metrics.percentiles_e2el_ms,
+        &config.selected_percentile_metrics,
+    );
+
+    // Steady-state metrics (null when not computed or scope gate failed).
+    result.insert(
+        "steady_state".into(),
+        serde_json::to_value(&metrics.steady_state).unwrap_or(Value::Null),
+    );
+
+    Value::Object(result)
+}
+
+/// Build multi-turn result JSON.
+pub fn build_multi_turn_result_json(
+    config: &BenchConfig,
+    mt_metrics: &MultiTurnMetrics,
+    conversation_outputs: &[ConversationOutput],
+    benchmark_duration: f64,
+    current_dt: &str,
+    spec_decode_stats: Option<&SpecDecodeStats>,
+) -> serde_json::Value {
+    let mut result = serde_json::Map::new();
+
+    // Setup
+    result.insert("date".into(), Value::String(current_dt.to_string()));
+    result.insert("mode".into(), Value::String("multi_turn".to_string()));
+    result.insert("backend".into(), Value::String(config.backend.to_string()));
+    result.insert(
+        "model_id".into(),
+        Value::String(config.model.clone().unwrap_or_default()),
+    );
+    result.insert(
+        "num_conversations".into(),
+        Value::Number(config.num_prompts.into()),
+    );
+    result.insert(
+        "max_model_len".into(),
+        config.max_model_len.map(|v| Value::Number(v.into())).unwrap_or(Value::Null),
+    );
+    result.insert(
+        "turns_per_conversation".into(),
+        Value::Number(config.multi_turn_num_turns.into()),
+    );
+    result.insert(
+        "multi_turn_concurrency".into(),
+        serde_json::json!(
+            config
+                .multi_turn_concurrency
+                .unwrap_or(config.max_concurrency.unwrap_or(config.num_prompts))
+        ),
+    );
+    result.insert(
+        "inter_turn_delay_ms".into(),
+        Value::Number(config.multi_turn_delay_ms.into()),
+    );
+    if config.multi_turn_prefix_global_ratio > 0.0
+        || config.multi_turn_prefix_conversation_ratio > 0.0
+    {
+        result.insert(
+            "prefix_global_ratio".into(),
+            serde_json::json!(config.multi_turn_prefix_global_ratio),
+        );
+        result.insert(
+            "prefix_conversation_ratio".into(),
+            serde_json::json!(config.multi_turn_prefix_conversation_ratio),
+        );
+        result.insert(
+            "prefix_unique_ratio".into(),
+            serde_json::json!(
+                (1.0 - (config.multi_turn_prefix_global_ratio
+                    + config.multi_turn_prefix_conversation_ratio))
+                    .max(0.0)
+            ),
+        );
+        result.insert("no_history_accumulation".into(), serde_json::json!(true));
+    }
+
+    // Metadata
+    if let Some(ref metadata) = config.metadata {
+        for (k, v) in metadata {
+            result.insert(k.clone(), Value::String(v.clone()));
+        }
+    }
+
+    // Conversation-level stats
+    result.insert("duration".into(), serde_json::json!(benchmark_duration));
+    result.insert(
+        "conversations_completed".into(),
+        serde_json::json!(mt_metrics.conversations_completed),
+    );
+    result.insert(
+        "conversations_failed".into(),
+        serde_json::json!(mt_metrics.conversations_failed),
+    );
+    result.insert(
+        "avg_turns_completed".into(),
+        serde_json::json!(mt_metrics.avg_turns_completed),
+    );
+    result.insert(
+        "avg_conversation_duration_ms".into(),
+        serde_json::json!(mt_metrics.avg_conversation_duration_ms),
+    );
+
+    // Overall metrics
+    let overall = &mt_metrics.overall;
+    result.insert("completed".into(), serde_json::json!(overall.completed));
+    result.insert("failed".into(), serde_json::json!(overall.failed));
+
+    // Per-failure log (always emitted when there are failures). Each entry is a single failed turn.
+    if overall.failed > 0 {
+        result.insert(
+            "failed_requests".into(),
+            Value::Array(collect_failed_turns(conversation_outputs)),
+        );
+    }
+    result.insert(
+        "total_input_tokens".into(),
+        serde_json::json!(overall.total_input),
+    );
+    result.insert(
+        "total_output_tokens".into(),
+        serde_json::json!(overall.total_output),
+    );
+    result.insert(
+        "request_throughput".into(),
+        serde_json::json!(overall.request_throughput),
+    );
+    result.insert(
+        "input_throughput".into(),
+        serde_json::json!(overall.input_throughput),
+    );
+    result.insert(
+        "output_throughput".into(),
+        serde_json::json!(overall.output_throughput),
+    );
+    result.insert(
+        "total_token_throughput".into(),
+        serde_json::json!(overall.total_token_throughput),
+    );
+
+    // Overall per-metric stats
+    add_metric_stats(
+        &mut result,
+        "ttft",
+        overall.mean_ttft_ms,
+        overall.median_ttft_ms,
+        overall.std_ttft_ms,
+        &overall.percentiles_ttft_ms,
+        &config.selected_percentile_metrics,
+    );
+    add_metric_stats(
+        &mut result,
+        "tpot",
+        overall.mean_tpot_ms,
+        overall.median_tpot_ms,
+        overall.std_tpot_ms,
+        &overall.percentiles_tpot_ms,
+        &config.selected_percentile_metrics,
+    );
+    add_metric_stats(
+        &mut result,
+        "itl",
+        overall.mean_itl_ms,
+        overall.median_itl_ms,
+        overall.std_itl_ms,
+        &overall.percentiles_itl_ms,
+        &config.selected_percentile_metrics,
+    );
+    add_metric_stats(
+        &mut result,
+        "e2el",
+        overall.mean_e2el_ms,
+        overall.median_e2el_ms,
+        overall.std_e2el_ms,
+        &overall.percentiles_e2el_ms,
+        &config.selected_percentile_metrics,
+    );
+
+    // Speculative decoding stats
+    if let Some(stats) = spec_decode_stats {
+        insert_spec_decode_stats(&mut result, stats);
+    }
+
+    // Per-turn metrics array
+    let per_turn_json: Vec<Value> = mt_metrics
+        .per_turn
+        .iter()
+        .enumerate()
+        .map(|(i, m)| {
+            let mut turn = serde_json::Map::new();
+            turn.insert("turn_index".into(), serde_json::json!(i));
+            turn.insert(
+                "num_samples".into(),
+                serde_json::json!(m.completed + m.failed),
+            );
+            turn.insert("completed".into(), serde_json::json!(m.completed));
+            turn.insert("failed".into(), serde_json::json!(m.failed));
+            turn.insert(
+                "total_input_tokens".into(),
+                serde_json::json!(m.total_input),
+            );
+            turn.insert(
+                "total_output_tokens".into(),
+                serde_json::json!(m.total_output),
+            );
+            turn.insert(
+                "request_throughput".into(),
+                serde_json::json!(m.request_throughput),
+            );
+            turn.insert(
+                "input_throughput".into(),
+                serde_json::json!(m.input_throughput),
+            );
+            turn.insert(
+                "output_throughput".into(),
+                serde_json::json!(m.output_throughput),
+            );
+            turn.insert(
+                "total_token_throughput".into(),
+                serde_json::json!(m.total_token_throughput),
+            );
+
+            add_metric_stats(
+                &mut turn,
+                "ttft",
+                m.mean_ttft_ms,
+                m.median_ttft_ms,
+                m.std_ttft_ms,
+                &m.percentiles_ttft_ms,
+                &config.selected_percentile_metrics,
+            );
+            add_metric_stats(
+                &mut turn,
+                "tpot",
+                m.mean_tpot_ms,
+                m.median_tpot_ms,
+                m.std_tpot_ms,
+                &m.percentiles_tpot_ms,
+                &config.selected_percentile_metrics,
+            );
+            add_metric_stats(
+                &mut turn,
+                "itl",
+                m.mean_itl_ms,
+                m.median_itl_ms,
+                m.std_itl_ms,
+                &m.percentiles_itl_ms,
+                &config.selected_percentile_metrics,
+            );
+            add_metric_stats(
+                &mut turn,
+                "e2el",
+                m.mean_e2el_ms,
+                m.median_e2el_ms,
+                m.std_e2el_ms,
+                &m.percentiles_e2el_ms,
+                &config.selected_percentile_metrics,
+            );
+
+            Value::Object(turn)
+        })
+        .collect();
+
+    result.insert("per_turn_metrics".into(), Value::Array(per_turn_json));
+
+    Value::Object(result)
+}
+
+fn collect_failed_requests(outputs: &[RequestFuncOutput]) -> Vec<Value> {
+    outputs
+        .iter()
+        .enumerate()
+        .filter(|(_, o)| !o.success)
+        .map(|(i, o)| {
+            serde_json::json!({
+                "index": i,
+                "error": o.error,
+                "prompt_len": o.prompt_len,
+                "start_time": o.start_time,
+                "latency": o.latency,
+                "ttft": o.ttft,
+                "output_tokens": o.output_tokens,
+                "itl": o.itl,
+            })
+        })
+        .collect()
+}
+
+fn collect_failed_turns(conversation_outputs: &[ConversationOutput]) -> Vec<Value> {
+    let mut failed = Vec::new();
+    for conv in conversation_outputs {
+        for turn in &conv.turns {
+            if !turn.request_output.success {
+                let o = &turn.request_output;
+                failed.push(serde_json::json!({
+                    "conversation_id": conv.conversation_id,
+                    "turn_index": turn.turn_index,
+                    "error": o.error,
+                    "prompt_len": o.prompt_len,
+                    "start_time": o.start_time,
+                    "latency": o.latency,
+                    "ttft": o.ttft,
+                    "output_tokens": o.output_tokens,
+                    "itl": o.itl,
+                }));
+            }
+        }
+    }
+    failed
+}
+
+fn insert_spec_decode_stats(result: &mut serde_json::Map<String, Value>, stats: &SpecDecodeStats) {
+    result.insert(
+        "spec_decode_acceptance_rate".into(),
+        serde_json::json!(stats.acceptance_rate),
+    );
+    result.insert(
+        "spec_decode_acceptance_length".into(),
+        serde_json::json!(stats.acceptance_length),
+    );
+    result.insert(
+        "spec_decode_num_drafts".into(),
+        serde_json::json!(stats.num_drafts),
+    );
+    result.insert(
+        "spec_decode_draft_tokens".into(),
+        serde_json::json!(stats.draft_tokens),
+    );
+    result.insert(
+        "spec_decode_accepted_tokens".into(),
+        serde_json::json!(stats.accepted_tokens),
+    );
+    result.insert(
+        "spec_decode_per_position_acceptance_rates".into(),
+        serde_json::json!(stats.per_position_acceptance_rates),
+    );
+}
+
+fn add_metric_stats(
+    result: &mut serde_json::Map<String, Value>,
+    name: &str,
+    mean: f64,
+    median: f64,
+    std: f64,
+    percentiles: &[(f64, f64)],
+    selected: &[String],
+) {
+    if !selected.iter().any(|s| s == name) {
+        return;
+    }
+    result.insert(format!("mean_{name}_ms"), serde_json::json!(mean));
+    result.insert(format!("median_{name}_ms"), serde_json::json!(median));
+    result.insert(format!("std_{name}_ms"), serde_json::json!(std));
+    for (p, value) in percentiles {
+        let p_str = if *p == p.floor() {
+            format!("{}", *p as i64)
+        } else {
+            format!("{p}")
+        };
+        result.insert(format!("p{p_str}_{name}_ms"), serde_json::json!(value));
+    }
+}
+
+/// Save result JSON to file (overwrite mode).
+pub fn save_result(json: &Value, file_path: &str) -> Result<()> {
+    let content = serde_json::to_string(json)?;
+    std::fs::write(file_path, content)?;
+    println!("Results saved to {file_path}");
+    Ok(())
+}
+
+/// Append result JSON to file (JSONL format, matching Python's --append-result).
+pub fn append_result(json: &Value, file_path: &str) -> Result<()> {
+    use std::io::Write;
+    let content = serde_json::to_string(json)?;
+    let mut file = std::fs::OpenOptions::new().create(true).append(true).open(file_path)?;
+    // If file is non-empty, prepend a newline
+    let meta = file.metadata()?;
+    if meta.len() > 0 {
+        file.write_all(b"\n")?;
+    }
+    file.write_all(content.as_bytes())?;
+    println!("Results appended to {file_path}");
+    Ok(())
+}
+
+/// Compute the result filename matching Python's logic.
+pub fn compute_result_filename(config: &BenchConfig, model_id: &str, current_dt: &str) -> String {
+    let base_model = model_id.split('/').next_back().unwrap_or(model_id);
+    let max_conc_str =
+        config.max_concurrency.map(|mc| format!("-concurrency{mc}")).unwrap_or_default();
+    let label = config.label.as_deref().unwrap_or_else(|| config.backend.as_str());
+
+    let file_name = if config.request_rate.is_infinite() {
+        format!("{label}-infqps{max_conc_str}-{base_model}-{current_dt}.json")
+    } else {
+        format!(
+            "{label}-{}qps{max_conc_str}-{base_model}-{current_dt}.json",
+            config.request_rate
+        )
+    };
+
+    if let Some(ref explicit) = config.result_filename {
+        if let Some(ref dir) = config.result_dir {
+            return format!("{dir}/{explicit}");
+        }
+        return explicit.clone();
+    }
+
+    if let Some(ref dir) = config.result_dir {
+        format!("{dir}/{file_name}")
+    } else {
+        file_name
+    }
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+
+    #[test]
+    fn test_insert_spec_decode_stats() {
+        let stats = SpecDecodeStats {
+            num_drafts: 1000,
+            draft_tokens: 3000,
+            accepted_tokens: 2400,
+            acceptance_rate: 80.0,
+            acceptance_length: 3.4,
+            per_position_acceptance_rates: vec![0.9, 0.8, 0.7],
+        };
+
+        let mut result = serde_json::Map::new();
+        insert_spec_decode_stats(&mut result, &stats);
+
+        assert_eq!(result["spec_decode_acceptance_rate"], 80.0);
+        assert_eq!(result["spec_decode_acceptance_length"], 3.4);
+        assert_eq!(result["spec_decode_num_drafts"], 1000);
+        assert_eq!(result["spec_decode_draft_tokens"], 3000);
+        assert_eq!(result["spec_decode_accepted_tokens"], 2400);
+        assert_eq!(
+            result["spec_decode_per_position_acceptance_rates"],
+            serde_json::json!([0.9, 0.8, 0.7])
+        );
+    }
+
+    #[test]
+    fn test_insert_spec_decode_stats_empty_positions() {
+        let stats = SpecDecodeStats {
+            num_drafts: 500,
+            draft_tokens: 1500,
+            accepted_tokens: 1200,
+            acceptance_rate: 80.0,
+            acceptance_length: 3.4,
+            per_position_acceptance_rates: vec![],
+        };
+
+        let mut result = serde_json::Map::new();
+        insert_spec_decode_stats(&mut result, &stats);
+
+        assert_eq!(
+            result["spec_decode_per_position_acceptance_rates"],
+            serde_json::json!([])
+        );
+        assert_eq!(result.len(), 6);
+    }
+
+    #[test]
+    fn test_collect_failed_requests_empty_when_all_success() {
+        let outputs = vec![
+            RequestFuncOutput {
+                success: true,
+                ..Default::default()
+            },
+            RequestFuncOutput {
+                success: true,
+                ..Default::default()
+            },
+        ];
+        assert!(collect_failed_requests(&outputs).is_empty());
+    }
+
+    #[test]
+    fn test_collect_failed_requests_preserves_index_and_fields() {
+        let outputs = vec![
+            RequestFuncOutput {
+                success: true,
+                ..Default::default()
+            },
+            // Pre-stream failure: no tokens received.
+            RequestFuncOutput {
+                success: false,
+                error: "connection reset".into(),
+                prompt_len: 128,
+                start_time: 1.5,
+                latency: 0.0,
+                ttft: 0.0,
+                output_tokens: 0,
+                itl: vec![],
+                ..Default::default()
+            },
+            RequestFuncOutput {
+                success: true,
+                ..Default::default()
+            },
+            // Partial failure: streamed 3 tokens before timing out.
+            RequestFuncOutput {
+                success: false,
+                error: "timeout".into(),
+                prompt_len: 64,
+                start_time: 2.25,
+                latency: 30.0,
+                ttft: 0.5,
+                output_tokens: 3,
+                itl: vec![0.1, 0.12],
+                ..Default::default()
+            },
+        ];
+
+        let failed = collect_failed_requests(&outputs);
+        assert_eq!(failed.len(), 2);
+
+        assert_eq!(failed[0]["index"], 1);
+        assert_eq!(failed[0]["error"], "connection reset");
+        assert_eq!(failed[0]["prompt_len"], 128);
+        assert_eq!(failed[0]["start_time"], 1.5);
+        assert_eq!(failed[0]["latency"], 0.0);
+        assert_eq!(failed[0]["ttft"], 0.0);
+        assert_eq!(failed[0]["output_tokens"], 0);
+        assert_eq!(failed[0]["itl"], serde_json::json!([]));
+
+        assert_eq!(failed[1]["index"], 3);
+        assert_eq!(failed[1]["error"], "timeout");
+        assert_eq!(failed[1]["latency"], 30.0);
+        assert_eq!(failed[1]["ttft"], 0.5);
+        assert_eq!(failed[1]["output_tokens"], 3);
+        assert_eq!(failed[1]["itl"], serde_json::json!([0.1, 0.12]));
+    }
+}
@@ -0,0 +1,5 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+pub mod console;
+pub mod json;
@@ -0,0 +1,171 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+use rand::SeedableRng;
+use rand::rngs::StdRng;
+use rand_distr::{Distribution, Gamma};
+
+use crate::cli::RampUpStrategy;
+use crate::config::RampUpConfig;
+
+/// Pre-computed request schedule: cumulative delays from start and per-request rates.
+pub struct RequestSchedule {
+    /// Cumulative absolute delay from start (seconds) for each request.
+    pub delays: Vec<f64>,
+    /// Instantaneous rate for each request (for ramp-up logging).
+    #[allow(dead_code)]
+    pub rates: Vec<f64>,
+}
+
+/// Compute the per-request rate, accounting for optional ramp-up.
+///
+/// Mirrors Python's `_get_current_request_rate()` from serve.py:218-240.
+fn get_current_request_rate(
+    ramp_up: Option<&RampUpConfig>,
+    request_index: usize,
+    total_requests: usize,
+    base_rate: f64,
+) -> f64 {
+    let config = match ramp_up {
+        Some(c) => c,
+        None => return base_rate,
+    };
+
+    let progress = request_index as f64 / (total_requests - 1).max(1) as f64;
+
+    match config.strategy {
+        RampUpStrategy::Linear => {
+            let increase = (config.end_rps - config.start_rps) * progress;
+            config.start_rps + increase
+        }
+        RampUpStrategy::Exponential => {
+            let ratio = config.end_rps / config.start_rps;
+            config.start_rps * ratio.powf(progress)
+        }
+    }
+}
+
+/// Compute the request schedule for all requests.
+///
+/// Ports the Python `get_request()` / `_generate_request_timestamps()` logic
+/// from serve.py:243-340.
+pub fn compute_schedule(
+    num_requests: usize,
+    request_rate: f64,
+    burstiness: f64,
+    seed: u64,
+    ramp_up: Option<&RampUpConfig>,
+) -> RequestSchedule {
+    assert!(burstiness > 0.0, "burstiness must be positive");
+    assert!(num_requests > 0, "must have at least one request");
+
+    let mut rng = StdRng::seed_from_u64(seed);
+    let mut delay_ts = Vec::with_capacity(num_requests);
+    let mut rates = Vec::with_capacity(num_requests);
+
+    for i in 0..num_requests {
+        let current_rate = get_current_request_rate(ramp_up, i, num_requests, request_rate);
+        rates.push(current_rate);
+
+        if current_rate.is_infinite() {
+            delay_ts.push(0.0);
+        } else if burstiness.is_infinite() {
+            // When burstiness → ∞, delay becomes constant = 1/rate
+            delay_ts.push(1.0 / current_rate);
+        } else {
+            let theta = 1.0 / (current_rate * burstiness);
+            let gamma = Gamma::new(burstiness, theta).unwrap();
+            delay_ts.push(gamma.sample(&mut rng));
+        }
+    }
+
+    // Compute cumulative delays
+    for i in 1..delay_ts.len() {
+        delay_ts[i] += delay_ts[i - 1];
+    }
+
+    // Normalize: scale cumulative delays so total matches target.
+    // Only for fixed-rate (no ramp-up) mode, matching Python behavior.
+    if ramp_up.is_none()
+        && let Some(&last) = delay_ts.last()
+        && last > 0.0
+        && !request_rate.is_infinite()
+    {
+        let target_total = num_requests as f64 / request_rate;
+        let factor = target_total / last;
+        for d in &mut delay_ts {
+            *d *= factor;
+        }
+    }
+
+    RequestSchedule {
+        delays: delay_ts,
+        rates,
+    }
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+
+    #[test]
+    fn test_infinite_rate_all_zero_delays() {
+        let sched = compute_schedule(100, f64::INFINITY, 1.0, 42, None);
+        assert_eq!(sched.delays.len(), 100);
+        for d in &sched.delays {
+            assert_eq!(*d, 0.0);
+        }
+    }
+
+    #[test]
+    fn test_fixed_rate_monotonic() {
+        let sched = compute_schedule(50, 10.0, 1.0, 42, None);
+        for i in 1..sched.delays.len() {
+            assert!(sched.delays[i] >= sched.delays[i - 1]);
+        }
+        // Total should be close to num_requests / rate = 5.0
+        let last = *sched.delays.last().unwrap();
+        assert!((last - 5.0).abs() < 0.01, "last delay = {last}");
+    }
+
+    #[test]
+    fn test_infinite_burstiness_constant_delay() {
+        let sched = compute_schedule(10, 5.0, f64::INFINITY, 42, None);
+        // Each delay should be 0.2s apart (1/5)
+        for i in 1..sched.delays.len() {
+            let delta = sched.delays[i] - sched.delays[i - 1];
+            assert!((delta - 0.2).abs() < 1e-10);
+        }
+    }
+
+    #[test]
+    fn test_linear_ramp_up() {
+        let ramp = RampUpConfig {
+            strategy: RampUpStrategy::Linear,
+            start_rps: 1.0,
+            end_rps: 10.0,
+        };
+        let sched = compute_schedule(10, 5.0, 1.0, 42, Some(&ramp));
+        // Rates should increase linearly from 1.0 to 10.0
+        assert!((sched.rates[0] - 1.0).abs() < 1e-10);
+        assert!((sched.rates[9] - 10.0).abs() < 1e-10);
+        // Middle should be ~5.5
+        assert!((sched.rates[4] - 5.0).abs() < 0.5);
+    }
+
+    #[test]
+    fn test_exponential_ramp_up() {
+        let ramp = RampUpConfig {
+            strategy: RampUpStrategy::Exponential,
+            start_rps: 1.0,
+            end_rps: 100.0,
+        };
+        let sched = compute_schedule(10, 5.0, 1.0, 42, Some(&ramp));
+        assert!((sched.rates[0] - 1.0).abs() < 1e-10);
+        assert!((sched.rates[9] - 100.0).abs() < 1e-10);
+        // Exponential: rates should be monotonically increasing
+        for i in 1..sched.rates.len() {
+            assert!(sched.rates[i] >= sched.rates[i - 1]);
+        }
+    }
+}
@@ -0,0 +1,72 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+use std::time::Instant;
+
+use indicatif::{ProgressBar, ProgressStyle};
+
+use crate::backends::{RequestFuncInput, RequestFuncOutput, get_backend};
+use crate::cli::BackendKind;
+use crate::error::{BenchError, Result};
+
+/// Wait for the serving endpoint to become available.
+///
+/// Sends test requests with retry until success or timeout.
+/// Mirrors Python's `wait_for_endpoint` in ready_checker.py.
+pub async fn wait_for_endpoint(
+    backend: BackendKind,
+    client: &reqwest::Client,
+    test_input: &RequestFuncInput,
+    timeout_seconds: u64,
+    retry_interval: u64,
+) -> Result<RequestFuncOutput> {
+    let backend = get_backend(backend)?;
+    let deadline = Instant::now() + std::time::Duration::from_secs(timeout_seconds);
+
+    println!("Waiting for endpoint to become up in {timeout_seconds}s");
+
+    let pb = ProgressBar::new(timeout_seconds);
+    pb.set_style(
+        ProgressStyle::with_template("{msg} |{bar:40}| {elapsed} elapsed, {eta} remaining")
+            .unwrap()
+            .progress_chars("##-"),
+    );
+
+    let mut last_error = String::new();
+
+    loop {
+        let remaining = deadline.saturating_duration_since(Instant::now());
+        let elapsed = timeout_seconds.saturating_sub(remaining.as_secs());
+        pb.set_position(elapsed);
+
+        if remaining.is_zero() {
+            pb.finish_and_clear();
+            break;
+        }
+
+        // Ping the endpoint
+        match backend.send_request(test_input, client).await {
+            Ok(output) if output.success => {
+                pb.finish_and_clear();
+                return Ok(output);
+            }
+            Ok(output) => {
+                let err = output.error.clone();
+                let err_last_line = err.lines().last().unwrap_or(&err);
+                eprintln!("Endpoint is not ready. Error='{err_last_line}'");
+                last_error = err;
+            }
+            Err(e) => {
+                last_error = e.to_string();
+            }
+        }
+
+        // Retry after delay
+        let sleep_dur = std::cmp::min(std::time::Duration::from_secs(retry_interval), remaining);
+        if !sleep_dur.is_zero() {
+            tokio::time::sleep(sleep_dur).await;
+        }
+    }
+
+    Err(BenchError::EndpointTimeout(timeout_seconds, last_error))
+}
@@ -0,0 +1,561 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+use crate::config::BenchConfig;
+use crate::error::{BenchError, Result};
+use crate::output::json::{compute_result_filename, save_result};
+
+/// Reset the server's prefix cache by calling POST /reset_prefix_cache.
+/// Requires VLLM_SERVER_DEV_MODE=1 on the vLLM server.
+async fn reset_prefix_cache(base_url: &str) -> Result<()> {
+    let url = format!("{}/reset_prefix_cache", base_url.trim_end_matches('/'));
+    let client = reqwest::Client::new();
+    let resp = client
+        .post(&url)
+        .send()
+        .await
+        .map_err(|e| BenchError::Backend(format!("Failed to reset prefix cache: {e}")))?;
+    if resp.status().is_success() {
+        println!("Prefix cache reset successfully.");
+    } else {
+        let status = resp.status();
+        let body = resp.text().await.unwrap_or_default();
+        return Err(BenchError::Backend(format!(
+            "Failed to reset prefix cache: HTTP {status} — {body}. \
+             Is VLLM_SERVER_DEV_MODE=1 set on the server?"
+        )));
+    }
+    Ok(())
+}
+
+/// Result of a single sweep point.
+struct SweepPoint {
+    label: String,
+    #[allow(dead_code)]
+    value: f64,
+    result_json: serde_json::Value,
+    ss_request_throughput: Option<f64>,
+    ss_output_throughput: Option<f64>,
+}
+
+/// Run a sweep over max-concurrency values.
+pub async fn run_concurrency_sweep(
+    base_config: &BenchConfig,
+    values: &[usize],
+    num_prompts_factor: Option<usize>,
+) -> Result<()> {
+    println!("{:=^70}", " Concurrency Sweep ");
+    println!(
+        "Sweeping --max-concurrency over {} values: {:?}",
+        values.len(),
+        values
+    );
+    println!();
+
+    let mut points = Vec::with_capacity(values.len());
+
+    let current_dt = chrono::Local::now().format("%Y%m%d-%H%M%S").to_string();
+
+    for (i, &mc) in values.iter().enumerate() {
+        println!(
+            "{:-^70}",
+            format!(" Run {}/{}: max_concurrency={} ", i + 1, values.len(), mc)
+        );
+
+        if base_config.reset_prefix_cache {
+            reset_prefix_cache(&base_config.base_url).await?;
+        }
+
+        let mut config = base_config.clone();
+        config.max_concurrency = Some(mc);
+        if let Some(factor) = num_prompts_factor {
+            config.num_prompts = mc * factor;
+        }
+        // Suppress per-run save (we save with concurrency suffix below)
+        config.save_result = false;
+        config.append_result = false;
+
+        let result = crate::benchmark::run_benchmark(&config).await?;
+
+        if base_config.save_result {
+            let model_id = config.model.as_deref().unwrap_or("unknown");
+            let file_name = compute_result_filename(&config, model_id, &current_dt);
+            if let Some(ref dir) = config.result_dir {
+                std::fs::create_dir_all(dir)?;
+            }
+            save_result(&result, &file_name)?;
+        }
+
+        points.push(SweepPoint {
+            label: format!("concurrency={mc}"),
+            value: mc as f64,
+            ss_request_throughput: ss_f64(&result, "request_throughput"),
+            ss_output_throughput: ss_f64(&result, "output_throughput"),
+            result_json: result,
+        });
+
+        println!();
+    }
+
+    print_sweep_summary(
+        "Max Concurrency",
+        &points,
+        &base_config.sweep_summary_percentiles,
+    );
+    Ok(())
+}
+
+/// Run a sweep over multi-turn concurrency values.
+pub async fn run_multi_turn_concurrency_sweep(
+    base_config: &BenchConfig,
+    values: &[usize],
+    num_prompts_factor: Option<usize>,
+) -> Result<()> {
+    println!("{:=^70}", " Multi-Turn Concurrency Sweep ");
+    println!(
+        "Sweeping --multi-turn-concurrency over {} values: {:?}",
+        values.len(),
+        values
+    );
+    println!();
+
+    let mut points = Vec::with_capacity(values.len());
+    let current_dt = chrono::Local::now().format("%Y%m%d-%H%M%S").to_string();
+
+    for (i, &mc) in values.iter().enumerate() {
+        println!(
+            "{:-^70}",
+            format!(
+                " Run {}/{}: multi_turn_concurrency={} ",
+                i + 1,
+                values.len(),
+                mc
+            )
+        );
+
+        if base_config.reset_prefix_cache {
+            reset_prefix_cache(&base_config.base_url).await?;
+        }
+
+        let mut config = base_config.clone();
+        config.multi_turn_concurrency = Some(mc);
+        // Set max_concurrency so compute_result_filename includes the
+        // concurrency suffix — without this every sweep point overwrites
+        // the same file.
+        config.max_concurrency = Some(mc);
+        if let Some(factor) = num_prompts_factor {
+            config.num_prompts = mc * factor;
+        }
+        // Suppress per-run save (we save with concurrency suffix below)
+        config.save_result = false;
+        config.append_result = false;
+
+        let result = crate::multi_turn::run_multi_turn_benchmark(&config).await?;
+
+        if base_config.save_result {
+            let model_id = config.model.as_deref().unwrap_or("unknown");
+            let file_name = compute_result_filename(&config, model_id, &current_dt);
+            if let Some(ref dir) = config.result_dir {
+                std::fs::create_dir_all(dir)?;
+            }
+            save_result(&result, &file_name)?;
+        }
+
+        points.push(SweepPoint {
+            label: format!("concurrency={mc}"),
+            value: mc as f64,
+            ss_request_throughput: ss_f64(&result, "request_throughput"),
+            ss_output_throughput: ss_f64(&result, "output_throughput"),
+            result_json: result,
+        });
+
+        println!();
+    }
+
+    print_sweep_summary(
+        "MT Concurrency",
+        &points,
+        &base_config.sweep_summary_percentiles,
+    );
+    Ok(())
+}
+
+/// Run a sweep over request-rate values.
+pub async fn run_rate_sweep(base_config: &BenchConfig, values: &[f64]) -> Result<()> {
+    println!("{:=^70}", " Request Rate Sweep ");
+    println!(
+        "Sweeping --request-rate over {} values: {:?}",
+        values.len(),
+        values
+    );
+    println!();
+
+    let mut points = Vec::with_capacity(values.len());
+    let current_dt = chrono::Local::now().format("%Y%m%d-%H%M%S").to_string();
+
+    for (i, &rate) in values.iter().enumerate() {
+        let rate_str = if rate.is_infinite() {
+            "inf".to_string()
+        } else {
+            format!("{rate}")
+        };
+        println!(
+            "{:-^70}",
+            format!(
+                " Run {}/{}: request_rate={} ",
+                i + 1,
+                values.len(),
+                rate_str
+            )
+        );
+
+        if base_config.reset_prefix_cache {
+            reset_prefix_cache(&base_config.base_url).await?;
+        }
+
+        let mut config = base_config.clone();
+        config.request_rate = rate;
+        config.save_result = false;
+        config.append_result = false;
+
+        let result = crate::benchmark::run_benchmark(&config).await?;
+
+        if base_config.save_result {
+            let model_id = config.model.as_deref().unwrap_or("unknown");
+            let file_name = compute_result_filename(&config, model_id, &current_dt);
+            if let Some(ref dir) = config.result_dir {
+                std::fs::create_dir_all(dir)?;
+            }
+            save_result(&result, &file_name)?;
+        }
+
+        points.push(SweepPoint {
+            label: format!("rate={rate_str}"),
+            value: rate,
+            ss_request_throughput: ss_f64(&result, "request_throughput"),
+            ss_output_throughput: ss_f64(&result, "output_throughput"),
+            result_json: result,
+        });
+
+        println!();
+    }
+
+    print_sweep_summary(
+        "Request Rate",
+        &points,
+        &base_config.sweep_summary_percentiles,
+    );
+    Ok(())
+}
+
+/// Print a summary table after all sweep points complete.
+fn print_sweep_summary(param_name: &str, points: &[SweepPoint], summary_percentiles: &[f64]) {
+    let param_width = param_name.len().max(20);
+    let columns = build_summary_columns(summary_percentiles);
+    let total_width = param_width
+        + columns.iter().map(SummaryColumn::render_width).sum::<usize>()
+        + columns.len();
+
+    println!("{:=^width$}", " Sweep Summary ", width = total_width);
+
+    print!("{param_name:<param_width$}");
+    for column in &columns {
+        print!(" {:>width$}", column.header, width = column.render_width());
+    }
+    println!();
+
+    print!("{:-<param_width$}", "");
+    for column in &columns {
+        print!(" {:-<width$}", "", width = column.render_width());
+    }
+    println!();
+
+    for point in points {
+        let row = render_summary_row(point, param_width, &columns);
+        println!("{row}");
+    }
+
+    println!("{:=<width$}", "", width = total_width);
+
+    // Find the best throughput point
+    if let Some(best) = points
+        .iter()
+        .filter(|p| get_f64(&p.result_json, "request_throughput").is_some())
+        .max_by(|a, b| {
+            let ta = get_f64(&a.result_json, "request_throughput").unwrap_or(0.0);
+            let tb = get_f64(&b.result_json, "request_throughput").unwrap_or(0.0);
+            ta.partial_cmp(&tb).unwrap_or(std::cmp::Ordering::Equal)
+        })
+    {
+        let tp = get_f64(&best.result_json, "request_throughput").unwrap_or(0.0);
+        println!("Best throughput: {:.2} req/s at {}", tp, best.label);
+    }
+}
+
+fn build_summary_columns(summary_percentiles: &[f64]) -> Vec<SummaryColumn> {
+    let mut columns = vec![
+        SummaryColumn::new("Req/s", "request_throughput", 10),
+        SummaryColumn::new("Tok/s", "output_throughput", 10),
+        SummaryColumn::new("Total tok/s", "total_token_throughput", 12),
+        SummaryColumn::new("SS req/s", SS_REQUEST_THROUGHPUT_KEY, 10),
+        SummaryColumn::new("SS out tok/s", SS_OUTPUT_THROUGHPUT_KEY, 12),
+        SummaryColumn::new("P50 TTFT(ms)", "median_ttft_ms", 12),
+        SummaryColumn::new("P50 TPOT(ms)", "median_tpot_ms", 12),
+        SummaryColumn::new("P90 TTFT(ms)", "p90_ttft_ms", 12),
+        SummaryColumn::new("P90 TPOT(ms)", "p90_tpot_ms", 12),
+    ];
+
+    for &percentile in summary_percentiles {
+        if percentile == 50.0 || percentile == 90.0 {
+            continue;
+        }
+        let p_str = format_percentile(percentile);
+        columns.push(SummaryColumn::new(
+            &format!("P{p_str} TTFT(ms)"),
+            &format!("p{p_str}_ttft_ms"),
+            12,
+        ));
+        columns.push(SummaryColumn::new(
+            &format!("P{p_str} TPOT(ms)"),
+            &format!("p{p_str}_tpot_ms"),
+            12,
+        ));
+    }
+
+    columns
+}
+
+fn render_summary_row(point: &SweepPoint, param_width: usize, columns: &[SummaryColumn]) -> String {
+    let mut row = format!("{:<param_width$}", point.label, param_width = param_width);
+    for column in columns {
+        let value = match column.key.as_str() {
+            SS_REQUEST_THROUGHPUT_KEY => point.ss_request_throughput,
+            SS_OUTPUT_THROUGHPUT_KEY => point.ss_output_throughput,
+            key => get_f64(&point.result_json, key),
+        };
+        row.push(' ');
+        row.push_str(&format!(
+            "{:>width$}",
+            fmt_f64(value),
+            width = column.render_width()
+        ));
+    }
+    row
+}
+
+fn format_percentile(percentile: f64) -> String {
+    if percentile == percentile.floor() {
+        format!("{}", percentile as i64)
+    } else {
+        format!("{percentile}")
+    }
+}
+
+struct SummaryColumn {
+    header: String,
+    key: String,
+    width: usize,
+}
+
+impl SummaryColumn {
+    fn new(header: &str, key: &str, width: usize) -> Self {
+        Self {
+            header: header.to_string(),
+            key: key.to_string(),
+            width,
+        }
+    }
+
+    fn render_width(&self) -> usize {
+        self.width.max(self.header.len())
+    }
+}
+
+/// Parse a comma-separated list of concurrency values.
+pub fn parse_concurrency_values(s: &str) -> Result<Vec<usize>> {
+    s.split(',')
+        .map(|v| {
+            v.trim().parse::<usize>().map_err(|_| {
+                BenchError::Config(format!("Invalid concurrency value: '{}'", v.trim()))
+            })
+        })
+        .collect()
+}
+
+/// Parse a comma-separated list of request rate values (supports "inf").
+pub fn parse_rate_values(s: &str) -> Result<Vec<f64>> {
+    s.split(',')
+        .map(|v| {
+            let v = v.trim();
+            if v == "inf" {
+                Ok(f64::INFINITY)
+            } else {
+                v.parse::<f64>()
+                    .map_err(|_| BenchError::Config(format!("Invalid request rate value: '{v}'")))
+            }
+        })
+        .collect()
+}
+
+fn get_f64(json: &serde_json::Value, key: &str) -> Option<f64> {
+    json.get(key).and_then(|v| v.as_f64())
+}
+
+/// Sentinel keys used on `SummaryColumn` to mark steady-state columns whose
+/// values come from `SweepPoint::ss_request_throughput` / `ss_output_throughput`
+/// rather than from the result JSON top level.
+const SS_REQUEST_THROUGHPUT_KEY: &str = "__ss_request_throughput";
+const SS_OUTPUT_THROUGHPUT_KEY: &str = "__ss_output_throughput";
+
+/// Extract a numeric field from the `steady_state` sub-object of a result JSON.
+/// Returns `None` when `steady_state` is missing or null, or when the field is
+/// absent / non-numeric.
+fn ss_f64(result: &serde_json::Value, key: &str) -> Option<f64> {
+    result
+        .get("steady_state")
+        .and_then(|ss| if ss.is_null() { None } else { Some(ss) })
+        .and_then(|ss| ss.get(key))
+        .and_then(|v| v.as_f64())
+}
+
+fn fmt_f64(v: Option<f64>) -> String {
+    match v {
+        Some(f) => format!("{:.2}", f),
+        None => "-".to_string(),
+    }
+}
+
+#[cfg(test)]
+mod tests {
+    use serde_json::json;
+
+    use super::*;
+
+    fn point(result_json: serde_json::Value) -> SweepPoint {
+        let ss_request_throughput = ss_f64(&result_json, "request_throughput");
+        let ss_output_throughput = ss_f64(&result_json, "output_throughput");
+        SweepPoint {
+            label: "point".to_string(),
+            value: 1.0,
+            result_json,
+            ss_request_throughput,
+            ss_output_throughput,
+        }
+    }
+
+    #[test]
+    fn summary_columns_follow_requested_percentile_order() {
+        let columns = build_summary_columns(&[90.0, 95.0]);
+        let headers: Vec<&str> = columns.iter().map(|column| column.header.as_str()).collect();
+
+        assert_eq!(
+            headers,
+            vec![
+                "Req/s",
+                "Tok/s",
+                "Total tok/s",
+                "SS req/s",
+                "SS out tok/s",
+                "P50 TTFT(ms)",
+                "P50 TPOT(ms)",
+                "P90 TTFT(ms)",
+                "P90 TPOT(ms)",
+                "P95 TTFT(ms)",
+                "P95 TPOT(ms)",
+            ]
+        );
+    }
+
+    #[test]
+    fn summary_row_uses_dash_for_missing_requested_percentile_values() {
+        let columns = build_summary_columns(&[90.0]);
+        let row = render_summary_row(
+            &point(json!({
+                "request_throughput": 1.0,
+                "output_throughput": 2.0,
+                "total_token_throughput": 3.0,
+                "median_ttft_ms": 10.0,
+                "median_tpot_ms": 20.0,
+                "p90_ttft_ms": 40.0,
+            })),
+            20,
+            &columns,
+        );
+
+        assert!(row.contains("40.00"));
+        assert!(row.contains(" -"));
+    }
+
+    #[test]
+    fn summary_columns_skip_duplicate_p50_and_p90_percentiles() {
+        let columns = build_summary_columns(&[50.0, 90.0]);
+        let headers: Vec<&str> = columns.iter().map(|column| column.header.as_str()).collect();
+
+        assert_eq!(
+            headers,
+            vec![
+                "Req/s",
+                "Tok/s",
+                "Total tok/s",
+                "SS req/s",
+                "SS out tok/s",
+                "P50 TTFT(ms)",
+                "P50 TPOT(ms)",
+                "P90 TTFT(ms)",
+                "P90 TPOT(ms)",
+            ]
+        );
+    }
+
+    #[test]
+    fn summary_row_renders_steady_state_columns() {
+        let columns = build_summary_columns(&[]);
+        let row = render_summary_row(
+            &point(json!({
+                "request_throughput": 1.0,
+                "output_throughput": 2.0,
+                "total_token_throughput": 3.0,
+                "median_ttft_ms": 10.0,
+                "median_tpot_ms": 20.0,
+                "p90_ttft_ms": 40.0,
+                "p90_tpot_ms": 50.0,
+                "steady_state": {
+                    "request_throughput": 0.77,
+                    "output_throughput": 88.88,
+                },
+            })),
+            20,
+            &columns,
+        );
+
+        assert!(row.contains("0.77"));
+        assert!(row.contains("88.88"));
+    }
+
+    #[test]
+    fn summary_row_renders_dash_when_steady_state_missing() {
+        let columns = build_summary_columns(&[]);
+        let row = render_summary_row(
+            &point(json!({
+                "request_throughput": 1.0,
+                "output_throughput": 2.0,
+                "total_token_throughput": 3.0,
+                "median_ttft_ms": 10.0,
+                "median_tpot_ms": 20.0,
+                "p90_ttft_ms": 40.0,
+                "p90_tpot_ms": 50.0,
+                "steady_state": null,
+            })),
+            20,
+            &columns,
+        );
+
+        // Two dashes for the two SS columns (SS req/s, SS out tok/s). Other
+        // columns have values so they won't emit dashes.
+        let dash_count = row.matches(" -").count();
+        assert!(
+            dash_count >= 2,
+            "expected at least 2 dashes, got {dash_count}: {row}"
+        );
+    }
+}
@@ -0,0 +1,542 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+use std::collections::HashMap;
+use std::path::Path;
+
+use rustc_hash::FxHashMap;
+
+use crate::error::{BenchError, Result};
+
+/// Default regex pattern matching cl100k_base (GPT-4, Qwen, etc.)
+const DEFAULT_TIKTOKEN_PATTERN: &str = r"(?i:'s|'t|'re|'ve|'m|'ll|'d)|[^\r\n\p{L}\p{N}]?\p{L}+|\p{N}{1,3}| ?[^\s\p{L}\p{N}]+[\r\n]*|\s*[\r\n]+|\s+(?!\S)|\s+";
+
+/// Number of reserved special token slots (matches Python's
+/// TikTokenTokenizer.num_reserved_special_tokens)
+const NUM_RESERVED_SPECIAL_TOKENS: u32 = 256;
+
+/// Tiktoken-based tokenizer for models that use tiktoken format (Kimi, Qwen, etc.)
+pub struct TiktokenTokenizer {
+    bpe: tiktoken_rs::CoreBPE,
+    vocab_size: u32,
+    #[allow(dead_code)]
+    num_base_tokens: u32,
+    /// IDs of all special tokens (for filtering from allowed_tokens)
+    special_token_ids: Vec<u32>,
+    /// Reverse mapping: token_id -> byte sequence, for lossy UTF-8 decoding.
+    /// Empty for built-in encodings (use bpe.decode instead).
+    decoder: Vec<Vec<u8>>,
+    /// True for built-in encodings loaded via tiktoken_rs (o200k_base, cl100k_base, etc.)
+    is_builtin: bool,
+}
+
+impl TiktokenTokenizer {
+    /// Load from a tiktoken .model file.
+    ///
+    /// `special_tokens_from_config`: tokens from tokenizer_config.json's added_tokens_decoder
+    /// `all_special_tokens`: ALL special tokens including the full 256 reserved slots
+    pub fn from_file(
+        model_path: &Path,
+        all_special_tokens: HashMap<String, u32>,
+        special_token_ids: Vec<u32>,
+        pattern: Option<&str>,
+    ) -> Result<Self> {
+        let content = std::fs::read_to_string(model_path)
+            .map_err(|e| BenchError::Tokenizer(format!("Failed to read tiktoken model: {e}")))?;
+
+        let mut encoder: FxHashMap<Vec<u8>, u32> = FxHashMap::default();
+
+        for line in content.lines() {
+            let line = line.trim();
+            if line.is_empty() {
+                continue;
+            }
+            let mut parts = line.split_whitespace();
+            let b64 = match parts.next() {
+                Some(s) => s,
+                None => continue,
+            };
+            let rank_str = match parts.next() {
+                Some(s) => s,
+                None => continue,
+            };
+
+            use base64::Engine;
+            let token_bytes = base64::engine::general_purpose::STANDARD
+                .decode(b64)
+                .map_err(|e| BenchError::Tokenizer(format!("Invalid base64 in model file: {e}")))?;
+            let rank: u32 = rank_str
+                .parse()
+                .map_err(|e| BenchError::Tokenizer(format!("Invalid rank: {e}")))?;
+            encoder.insert(token_bytes, rank);
+        }
+
+        if encoder.is_empty() {
+            return Err(BenchError::Tokenizer("Empty tiktoken model file".into()));
+        }
+
+        let num_base_tokens = encoder.len() as u32;
+
+        // Convert special tokens to FxHashMap for tiktoken-rs
+        let special_fx: FxHashMap<String, u32> =
+            all_special_tokens.iter().map(|(k, &v)| (k.clone(), v)).collect();
+
+        // vocab_size = base tokens + all reserved special token slots
+        let vocab_size = num_base_tokens + NUM_RESERVED_SPECIAL_TOKENS;
+
+        // Build reverse mapping for lossy decode
+        let mut decoder = vec![Vec::new(); vocab_size as usize];
+        for (bytes, &rank) in &encoder {
+            if (rank as usize) < decoder.len() {
+                decoder[rank as usize] = bytes.clone();
+            }
+        }
+        for (text, &rank) in &all_special_tokens {
+            if (rank as usize) < decoder.len() {
+                decoder[rank as usize] = text.as_bytes().to_vec();
+            }
+        }
+
+        let pat = pattern.unwrap_or(DEFAULT_TIKTOKEN_PATTERN);
+
+        let bpe = tiktoken_rs::CoreBPE::new(encoder, special_fx, pat)
+            .map_err(|e| BenchError::Tokenizer(format!("Failed to build tiktoken BPE: {e}")))?;
+
+        Ok(Self {
+            bpe,
+            vocab_size,
+            num_base_tokens,
+            special_token_ids,
+            decoder,
+            is_builtin: false,
+        })
+    }
+
+    /// Create from a built-in tiktoken encoding (o200k_base, cl100k_base, etc.)
+    pub fn from_builtin_bpe(bpe: tiktoken_rs::CoreBPE, vocab_size: u32) -> Self {
+        Self {
+            bpe,
+            vocab_size,
+            num_base_tokens: vocab_size,
+            special_token_ids: Vec::new(),
+            decoder: Vec::new(),
+            is_builtin: true,
+        }
+    }
+
+    pub fn encode(&self, text: &str) -> Vec<u32> {
+        // Use encode_with_special_tokens to match Python's encode(allowed_special="all")
+        self.bpe.encode_with_special_tokens(text)
+    }
+
+    /// Decode token IDs to text with lossy UTF-8 handling.
+    pub fn decode(&self, ids: &[u32]) -> Result<String> {
+        if self.is_builtin {
+            // Byte-level decode + lossy UTF-8, matching the file-based path below.
+            // (CoreBPE::decode errors on invalid UTF-8, which random token
+            // sequences routinely produce with byte-level BPE vocabularies.)
+            let bytes: Vec<u8> =
+                self.bpe._decode_native_and_split(ids.to_vec()).flatten().collect();
+            return Ok(String::from_utf8_lossy(&bytes).into_owned());
+        }
+        let mut bytes = Vec::new();
+        for &id in ids {
+            if let Some(token_bytes) = self.decoder.get(id as usize) {
+                bytes.extend_from_slice(token_bytes);
+            }
+        }
+        Ok(String::from_utf8_lossy(&bytes).into_owned())
+    }
+
+    pub fn vocab_size(&self) -> u32 {
+        self.vocab_size
+    }
+
+    /// Get non-special token IDs whose byte representation is valid UTF-8.
+    /// Excludes ALL special tokens (like Python's `set(all_tokens) - set(prohibited_tokens)`).
+    pub fn get_allowed_tokens(&self) -> Vec<u32> {
+        if self.is_builtin {
+            // For built-in encodings, return full token range.
+            // Note: when used with random dataset, these IDs are sent to vLLM as-is;
+            // ensure the model's tokenizer is compatible (e.g. GPT-4o for o200k_base).
+            return (0..self.vocab_size).collect();
+        }
+        let special_set: std::collections::HashSet<u32> =
+            self.special_token_ids.iter().copied().collect();
+
+        self.decoder
+            .iter()
+            .enumerate()
+            .filter(|(id, bytes)| {
+                !bytes.is_empty()
+                    && std::str::from_utf8(bytes).is_ok()
+                    && !special_set.contains(&(*id as u32))
+            })
+            .map(|(id, _)| id as u32)
+            .collect()
+    }
+}
+
+/// Load a built-in tiktoken encoding by name.
+///
+/// Supported names: `o200k_base` (GPT-4o), `cl100k_base` (GPT-4/3.5-turbo),
+/// `p50k_base`, `r50k_base`, `gpt2`.
+///
+/// These encodings are bundled with tiktoken-rs — no network download required.
+/// Useful for consistent cross-model token counting (e.g. Artificial Analysis methodology).
+pub fn load_builtin_tiktoken(encoding: &str) -> Result<TiktokenTokenizer> {
+    let (bpe, vocab_size) = match encoding {
+        "o200k_base" => (tiktoken_rs::o200k_base(), 200_275u32),
+        "cl100k_base" => (tiktoken_rs::cl100k_base(), 100_277u32),
+        "p50k_base" => (tiktoken_rs::p50k_base(), 50_281u32),
+        "p50k_edit" => (tiktoken_rs::p50k_edit(), 50_281u32),
+        "r50k_base" | "gpt2" => (tiktoken_rs::r50k_base(), 50_257u32),
+        _ => {
+            return Err(BenchError::Tokenizer(format!(
+                "Unknown built-in tiktoken encoding: '{encoding}'. \
+                 Supported: o200k_base, cl100k_base, p50k_base, r50k_base, gpt2"
+            )));
+        }
+    };
+    let bpe = bpe.map_err(|e| BenchError::Tokenizer(format!("Failed to load {encoding}: {e}")))?;
+    println!("Tokenizer: Built-in tiktoken {encoding} (vocab_size={vocab_size})");
+    Ok(TiktokenTokenizer::from_builtin_bpe(bpe, vocab_size))
+}
+
+/// Try to load a tiktoken tokenizer from a local directory or HuggingFace model repo.
+pub fn try_load_tiktoken(model_id: &str) -> Result<TiktokenTokenizer> {
+    // Phase 1: If model_id is a local directory, look for tiktoken files there
+    let local_dir = Path::new(model_id);
+    if local_dir.is_dir() {
+        return try_load_tiktoken_from_dir(local_dir, model_id);
+    }
+
+    // Phase 2: Fall back to HuggingFace Hub download
+    try_load_tiktoken_from_hf(model_id)
+}
+
+/// Common tiktoken model filenames to search for.
+const TIKTOKEN_MODEL_FILENAMES: &[&str] = &["tiktoken.model", "qwen.tiktoken", "vocab.tiktoken"];
+
+/// Load a tiktoken tokenizer from a local directory.
+fn try_load_tiktoken_from_dir(dir: &Path, model_id: &str) -> Result<TiktokenTokenizer> {
+    let model_path = TIKTOKEN_MODEL_FILENAMES
+        .iter()
+        .map(|f| dir.join(f))
+        .find(|p| p.exists())
+        .ok_or_else(|| {
+            BenchError::Tokenizer(format!(
+                "No tiktoken model file found in local directory '{model_id}' \
+                 (looked for: {})",
+                TIKTOKEN_MODEL_FILENAMES.join(", ")
+            ))
+        })?;
+
+    let num_base_tokens = count_base_tokens(&model_path)?;
+
+    let config_path = dir.join("tokenizer_config.json");
+    let config = if config_path.exists() {
+        read_tokenizer_config(&config_path)
+    } else {
+        None
+    };
+
+    let pattern = extract_pat_str_from_local_dir(dir);
+
+    build_tiktoken(model_id, &model_path, config, pattern, num_base_tokens)
+}
+
+/// Load a tiktoken tokenizer from a HuggingFace model repo.
+fn try_load_tiktoken_from_hf(model_id: &str) -> Result<TiktokenTokenizer> {
+    let repo = crate::hub::HubRepo::model(model_id.to_string());
+
+    let model_path = repo
+        .get("tiktoken.model")
+        .or_else(|_| repo.get("qwen.tiktoken"))
+        .or_else(|_| repo.get("vocab.tiktoken"))
+        .map_err(|_| {
+            BenchError::Tokenizer(format!("No tiktoken model file found for '{model_id}'"))
+        })?;
+
+    let num_base_tokens = count_base_tokens(&model_path)?;
+
+    let config = match repo.get("tokenizer_config.json") {
+        Ok(config_path) => read_tokenizer_config(&config_path),
+        Err(_) => None,
+    };
+
+    let pattern = extract_pat_str_from_repo(&repo);
+
+    build_tiktoken(model_id, &model_path, config, pattern, num_base_tokens)
+}
+
+/// Build a TiktokenTokenizer from discovered model file, config, and pattern.
+fn build_tiktoken(
+    model_id: &str,
+    model_path: &Path,
+    config: Option<TokenizerConfig>,
+    pattern: Option<String>,
+    num_base_tokens: u32,
+) -> Result<TiktokenTokenizer> {
+    // Build the full 256 reserved special tokens map.
+    // Python: {special_tokens_mapping.get(i, f"<|reserved_token_{i}|>"): i
+    //          for i in range(num_base_tokens, num_base_tokens + 256)}
+    let config_special = config.as_ref().map(|c| &c.added_tokens).cloned().unwrap_or_default();
+
+    // Invert config_special: id -> content
+    let id_to_content: HashMap<u32, String> =
+        config_special.iter().map(|(content, &id)| (id, content.clone())).collect();
+
+    let mut all_special_tokens: HashMap<String, u32> = HashMap::new();
+    let mut special_token_ids: Vec<u32> = Vec::new();
+
+    for i in 0..NUM_RESERVED_SPECIAL_TOKENS {
+        let token_id = num_base_tokens + i;
+        let content = id_to_content
+            .get(&token_id)
+            .cloned()
+            .unwrap_or_else(|| format!("<|reserved_token_{i}|>"));
+        all_special_tokens.insert(content, token_id);
+        special_token_ids.push(token_id);
+    }
+
+    // Also collect special IDs that Python marks as prohibited
+    // (those with "special": true in added_tokens_decoder)
+    let prohibited_ids = config.as_ref().map(|c| c.special_ids.clone()).unwrap_or_default();
+    for id in &prohibited_ids {
+        if !special_token_ids.contains(id) {
+            special_token_ids.push(*id);
+        }
+    }
+
+    println!(
+        "Loading tiktoken model for '{model_id}' (base={}, special={}, pat={})...",
+        num_base_tokens,
+        all_special_tokens.len(),
+        if pattern.is_some() {
+            "custom"
+        } else {
+            "default"
+        },
+    );
+
+    TiktokenTokenizer::from_file(
+        model_path,
+        all_special_tokens,
+        special_token_ids,
+        pattern.as_deref(),
+    )
+}
+
+/// Count base tokens in a tiktoken .model file without fully parsing it.
+fn count_base_tokens(model_path: &Path) -> Result<u32> {
+    let content = std::fs::read_to_string(model_path)
+        .map_err(|e| BenchError::Tokenizer(format!("Failed to read tiktoken model: {e}")))?;
+    let count = content
+        .lines()
+        .filter(|l| {
+            let l = l.trim();
+            !l.is_empty() && l.split_whitespace().count() >= 2
+        })
+        .count();
+    Ok(count as u32)
+}
+
+/// Parsed tokenizer config
+struct TokenizerConfig {
+    /// All tokens from added_tokens_decoder (content -> id)
+    added_tokens: HashMap<String, u32>,
+    /// IDs of tokens marked "special": true
+    special_ids: Vec<u32>,
+}
+
+/// Read and parse tokenizer_config.json
+fn read_tokenizer_config(config_path: &Path) -> Option<TokenizerConfig> {
+    let content = std::fs::read_to_string(config_path).ok()?;
+    let config: serde_json::Value = serde_json::from_str(&content).ok()?;
+
+    let mut added_tokens = HashMap::new();
+    let mut special_ids = Vec::new();
+
+    if let Some(added) = config.get("added_tokens_decoder").and_then(|v| v.as_object()) {
+        for (id_str, token_info) in added {
+            if let (Ok(id), Some(content)) = (
+                id_str.parse::<u32>(),
+                token_info.get("content").and_then(|c| c.as_str()),
+            ) {
+                added_tokens.insert(content.to_string(), id);
+
+                let is_special =
+                    token_info.get("special").and_then(|s| s.as_bool()).unwrap_or(false);
+                if is_special {
+                    special_ids.push(id);
+                }
+            }
+        }
+    }
+
+    Some(TokenizerConfig {
+        added_tokens,
+        special_ids,
+    })
+}
+
+/// Try to extract pat_str from Python tokenizer source files in a local directory.
+/// Returns None if unavailable or unparsable.
+fn extract_pat_str_from_local_dir(dir: &Path) -> Option<String> {
+    ["tokenization_kimi.py", "tokenizer.py"]
+        .iter()
+        .map(|f| dir.join(f))
+        .filter(|p| p.exists())
+        .find_map(|p| {
+            std::fs::read_to_string(p)
+                .ok()
+                .and_then(|source| extract_pat_str_from_source(&source))
+        })
+}
+
+/// Try to download the Python tokenizer source file and extract pat_str via regex.
+/// Returns None if unavailable or unparsable.
+fn extract_pat_str_from_repo(repo: &crate::hub::HubRepo) -> Option<String> {
+    // Try common Python tokenizer filenames
+    let py_path = repo.get("tokenization_kimi.py").or_else(|_| repo.get("tokenizer.py")).ok()?;
+
+    let source = std::fs::read_to_string(&py_path).ok()?;
+
+    // Look for pat_str assignment. Common patterns:
+    // pat_str = "..." or pat_str = '...' or pat_str = "|".join([...])
+    extract_pat_str_from_source(&source)
+}
+
+/// Parse pat_str from Python source code.
+fn extract_pat_str_from_source(source: &str) -> Option<String> {
+    // Strategy: find `pat_str = "|".join([` and collect the raw string fragments
+    // This handles the common Kimi/Qwen pattern of joining a list of regex strings.
+
+    // First try: look for pat_str = "|".join([...]) pattern
+    if let Some(join_start) = source.find("pat_str") {
+        let after = &source[join_start..];
+
+        // Check for "|".join([ pattern
+        if let Some(join_pos) = after.find(".join(") {
+            let after_join = &after[join_pos + 6..]; // skip ".join("
+            if let Some(bracket_start) = after_join.find('[') {
+                let inside = &after_join[bracket_start + 1..];
+                // Collect all string literals inside the list
+                let mut fragments = Vec::new();
+                let mut remaining = inside;
+
+                while let Some(frag) = extract_next_python_string(remaining) {
+                    fragments.push(frag.0);
+                    remaining = frag.1;
+                    // Check if we hit the closing bracket
+                    let trimmed = remaining.trim_start();
+                    if trimmed.starts_with(']') {
+                        break;
+                    }
+                }
+
+                if !fragments.is_empty() {
+                    let pattern = fragments.join("|");
+                    println!(
+                        "Extracted pat_str from Python source: {} fragments",
+                        fragments.len()
+                    );
+                    return Some(pattern);
+                }
+            }
+        }
+
+        // Fallback: simple pat_str = r"..." or pat_str = "..."
+        if let Some(eq_pos) = after.find('=') {
+            let after_eq = after[eq_pos + 1..].trim_start();
+            if let Some(frag) = extract_next_python_string(after_eq) {
+                return Some(frag.0);
+            }
+        }
+    }
+
+    None
+}
+
+/// Extract the next Python string literal (r"...", "...", r'''...''', etc.)
+/// Returns (string_content, remaining_text)
+fn extract_next_python_string(s: &str) -> Option<(String, &str)> {
+    let s = s.trim_start_matches(|c: char| c == ',' || c.is_whitespace());
+
+    // Skip comments
+    if s.starts_with('#') {
+        let next_line = s.find('\n').map(|i| i + 1).unwrap_or(s.len());
+        return extract_next_python_string(&s[next_line..]);
+    }
+
+    // Check for r""" (triple-quoted raw string)
+    for prefix in &["r\"\"\"", "r'''"] {
+        if let Some(inner) = s.strip_prefix(prefix) {
+            let delim = &prefix[1..]; // """ or '''
+            if let Some(end) = inner.find(delim) {
+                let content = &inner[..end];
+                let rest = &inner[end + delim.len()..];
+                return Some((content.to_string(), rest));
+            }
+        }
+    }
+
+    // Check for r"..." (raw string)
+    if let Some(inner) = s.strip_prefix("r\"")
+        && let Some(end) = inner.find('"')
+    {
+        let content = &inner[..end];
+        let rest = &inner[end + 1..];
+        return Some((content.to_string(), rest));
+    }
+    if let Some(inner) = s.strip_prefix("r'")
+        && let Some(end) = inner.find('\'')
+    {
+        let content = &inner[..end];
+        let rest = &inner[end + 1..];
+        return Some((content.to_string(), rest));
+    }
+
+    // Check for "..." or '...' (regular string — same as raw for regex patterns)
+    if let Some(inner) = s.strip_prefix('"')
+        && let Some(end) = find_unescaped(inner, '"')
+    {
+        let content = &inner[..end];
+        let rest = &inner[end + 1..];
+        return Some((content.to_string(), rest));
+    }
+    if let Some(inner) = s.strip_prefix('\'')
+        && let Some(end) = find_unescaped(inner, '\'')
+    {
+        let content = &inner[..end];
+        let rest = &inner[end + 1..];
+        return Some((content.to_string(), rest));
+    }
+
+    // Check for closing bracket — stop
+    if s.starts_with(']') {
+        return None;
+    }
+
+    None
+}
+
+/// Find position of `ch` that is not preceded by a backslash.
+fn find_unescaped(s: &str, ch: char) -> Option<usize> {
+    let mut escaped = false;
+    for (i, c) in s.char_indices() {
+        if escaped {
+            escaped = false;
+            continue;
+        }
+        if c == '\\' {
+            escaped = true;
+            continue;
+        }
+        if c == ch {
+            return Some(i);
+        }
+    }
+    None
+}
@@ -0,0 +1,300 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+use std::collections::HashSet;
+use std::path::Path;
+
+use tokenizers::Tokenizer;
+
+use crate::error::{BenchError, Result};
+use crate::tiktoken::TiktokenTokenizer;
+
+/// Abstraction over local HuggingFace tokenizer, tiktoken, or server-side tokenization.
+pub enum TokenizerKind {
+    Local(Box<Tokenizer>),
+    Tiktoken(TiktokenTokenizer),
+    Server(ServerTokenizer),
+}
+
+/// Server-side tokenizer using vLLM's /tokenize and /detokenize endpoints.
+pub struct ServerTokenizer {
+    client: reqwest::blocking::Client,
+    tokenize_url: String,
+    detokenize_url: String,
+    model: String,
+    cached_vocab_size: u32,
+}
+
+impl ServerTokenizer {
+    /// Create a new server tokenizer and verify connectivity.
+    pub fn new(base_url: &str, model: &str) -> Result<Self> {
+        let client = reqwest::blocking::Client::builder()
+            .timeout(std::time::Duration::from_secs(30))
+            .build()
+            .map_err(|e| BenchError::Tokenizer(format!("Failed to build HTTP client: {e}")))?;
+
+        let tokenize_url = format!("{base_url}/tokenize");
+        let detokenize_url = format!("{base_url}/detokenize");
+
+        let st = Self {
+            client,
+            tokenize_url,
+            detokenize_url,
+            model: model.to_string(),
+            cached_vocab_size: 0,
+        };
+
+        // Probe the endpoint to verify it works and discover vocab size
+        let test_tokens = st.encode_inner("test")?;
+        let max_id = test_tokens.iter().copied().max().unwrap_or(0);
+        let estimated_vocab = (max_id * 2).max(131072);
+
+        Ok(Self {
+            cached_vocab_size: estimated_vocab,
+            ..st
+        })
+    }
+
+    fn encode_inner(&self, text: &str) -> Result<Vec<u32>> {
+        let payload = serde_json::json!({
+            "model": self.model,
+            "prompt": text,
+        });
+
+        let resp = self
+            .client
+            .post(&self.tokenize_url)
+            .json(&payload)
+            .send()
+            .map_err(|e| BenchError::Tokenizer(format!("Server tokenize failed: {e}")))?;
+
+        if !resp.status().is_success() {
+            return Err(BenchError::Tokenizer(format!(
+                "Server tokenize returned HTTP {}",
+                resp.status()
+            )));
+        }
+
+        let data: serde_json::Value = resp.json().map_err(|e| {
+            BenchError::Tokenizer(format!("Failed to parse tokenize response: {e}"))
+        })?;
+
+        let tokens = data
+            .get("tokens")
+            .and_then(|t| t.as_array())
+            .ok_or_else(|| BenchError::Tokenizer("Missing 'tokens' in tokenize response".into()))?;
+
+        tokens
+            .iter()
+            .map(|v| {
+                v.as_u64()
+                    .map(|id| id as u32)
+                    .ok_or_else(|| BenchError::Tokenizer("Invalid token ID in response".into()))
+            })
+            .collect()
+    }
+
+    fn decode_inner(&self, ids: &[u32]) -> Result<String> {
+        let payload = serde_json::json!({
+            "model": self.model,
+            "tokens": ids,
+        });
+
+        let resp = self
+            .client
+            .post(&self.detokenize_url)
+            .json(&payload)
+            .send()
+            .map_err(|e| BenchError::Tokenizer(format!("Server detokenize failed: {e}")))?;
+
+        if !resp.status().is_success() {
+            return Err(BenchError::Tokenizer(format!(
+                "Server detokenize returned HTTP {}",
+                resp.status()
+            )));
+        }
+
+        let data: serde_json::Value = resp.json().map_err(|e| {
+            BenchError::Tokenizer(format!("Failed to parse detokenize response: {e}"))
+        })?;
+
+        data.get("prompt")
+            .and_then(|p| p.as_str())
+            .map(|s| s.to_string())
+            .ok_or_else(|| BenchError::Tokenizer("Missing 'prompt' in detokenize response".into()))
+    }
+}
+
+// --- TokenizerKind methods ---
+
+impl TokenizerKind {
+    pub fn encode(&self, text: &str, add_special_tokens: bool) -> Result<Vec<u32>> {
+        match self {
+            TokenizerKind::Local(tok) => tok
+                .encode(text, add_special_tokens)
+                .map(|enc| enc.get_ids().to_vec())
+                .map_err(|e| BenchError::Tokenizer(format!("Encode failed: {e}"))),
+            TokenizerKind::Tiktoken(tok) => Ok(tok.encode(text)),
+            TokenizerKind::Server(srv) => srv.encode_inner(text),
+        }
+    }
+
+    pub fn decode(&self, ids: &[u32], skip_special_tokens: bool) -> Result<String> {
+        match self {
+            TokenizerKind::Local(tok) => tok
+                .decode(ids, skip_special_tokens)
+                .map_err(|e| BenchError::Tokenizer(format!("Decode failed: {e}"))),
+            TokenizerKind::Tiktoken(tok) => tok.decode(ids),
+            TokenizerKind::Server(srv) => srv.decode_inner(ids),
+        }
+    }
+
+    pub fn vocab_size(&self) -> u32 {
+        match self {
+            TokenizerKind::Local(tok) => tok.get_vocab_size(true) as u32,
+            TokenizerKind::Tiktoken(tok) => tok.vocab_size(),
+            TokenizerKind::Server(srv) => srv.cached_vocab_size,
+        }
+    }
+
+    pub fn num_special_tokens_to_add(&self) -> usize {
+        match self {
+            TokenizerKind::Local(tok) => match tok.encode("", true) {
+                Ok(enc) => enc.get_ids().len(),
+                Err(_) => 0,
+            },
+            TokenizerKind::Tiktoken(_) | TokenizerKind::Server(_) => 0,
+        }
+    }
+
+    pub fn get_allowed_tokens(&self) -> Vec<u32> {
+        match self {
+            TokenizerKind::Local(tok) => {
+                let vs = tok.get_vocab_size(true) as u32;
+                let mut special_ids = HashSet::new();
+                for (id, token) in tok.get_added_tokens_decoder() {
+                    if token.special {
+                        special_ids.insert(id);
+                    }
+                }
+                (0..vs).filter(|id| !special_ids.contains(id)).collect()
+            }
+            TokenizerKind::Tiktoken(tok) => tok.get_allowed_tokens(),
+            TokenizerKind::Server(srv) => (0..srv.cached_vocab_size).collect(),
+        }
+    }
+}
+
+/// Load a tokenizer with fallback chain:
+/// 0. Built-in tiktoken encoding (o200k_base, cl100k_base, etc.) — no download needed
+/// 1. Local tokenizer.json (HuggingFace fast tokenizer)
+/// 2. Tiktoken model file (for Kimi, Qwen, etc.)
+/// 3. Server-side /tokenize + /detokenize endpoints
+///
+/// `server_info` is `Some((base_url, model))` to enable server-side fallback.
+pub fn load_tokenizer(
+    model_id: &str,
+    _trust_remote_code: bool,
+    server_info: Option<(&str, &str)>,
+) -> Result<TokenizerKind> {
+    // 0. Check for built-in tiktoken encoding names (no HF download needed). These are useful for
+    //    consistent cross-model token counting (e.g. Artificial Analysis).
+    const BUILTIN_TIKTOKEN: &[&str] = &[
+        "o200k_base",
+        "cl100k_base",
+        "p50k_base",
+        "p50k_edit",
+        "r50k_base",
+        "gpt2",
+    ];
+    if BUILTIN_TIKTOKEN.contains(&model_id) {
+        return crate::tiktoken::load_builtin_tiktoken(model_id).map(TokenizerKind::Tiktoken);
+    }
+
+    // 1. Try local HuggingFace tokenizer (tokenizer.json)
+    match try_load_local(model_id) {
+        Ok(tok) => {
+            println!("Tokenizer: Local (vocab_size={})", tok.get_vocab_size(true));
+            Ok(TokenizerKind::Local(Box::new(tok)))
+        }
+        Err(local_err) => {
+            // 2. Try tiktoken format
+            println!("No tokenizer.json for '{model_id}', trying tiktoken format...");
+            match crate::tiktoken::try_load_tiktoken(model_id) {
+                Ok(tok) => {
+                    println!("Tokenizer: Tiktoken (vocab_size={})", tok.vocab_size());
+                    Ok(TokenizerKind::Tiktoken(tok))
+                }
+                Err(tiktoken_err) => {
+                    // 3. Try server-side fallback
+                    if let Some((base_url, model)) = server_info {
+                        println!(
+                            "Tiktoken also not available ({tiktoken_err}), \
+                             trying server-side tokenization..."
+                        );
+                        match ServerTokenizer::new(base_url, model) {
+                            Ok(srv) => {
+                                println!(
+                                    "Tokenizer: Server (vocab_size≈{})",
+                                    srv.cached_vocab_size
+                                );
+                                return Ok(TokenizerKind::Server(srv));
+                            }
+                            Err(srv_err) => {
+                                return Err(BenchError::Tokenizer(format!(
+                                    "All tokenizer loading methods failed:\n  \
+                                     Local: {local_err}\n  \
+                                     Tiktoken: {tiktoken_err}\n  \
+                                     Server: {srv_err}\n  \
+                                     Try --tokenizer with a model that has tokenizer.json."
+                                )));
+                            }
+                        }
+                    }
+                    Err(BenchError::Tokenizer(format!(
+                        "Failed to load tokenizer:\n  \
+                         Local: {local_err}\n  \
+                         Tiktoken: {tiktoken_err}\n  \
+                         Try --tokenizer or provide --base-url for server fallback."
+                    )))
+                }
+            }
+        }
+    }
+}
+
+/// Try loading tokenizer.json from local path or HuggingFace Hub.
+fn try_load_local(model_id: &str) -> Result<Tokenizer> {
+    // 1. Try local directory with tokenizer.json
+    let local_path = Path::new(model_id).join("tokenizer.json");
+    if local_path.exists() {
+        return Tokenizer::from_file(&local_path).map_err(|e| {
+            BenchError::Tokenizer(format!(
+                "Failed to load tokenizer from {}: {e}",
+                local_path.display()
+            ))
+        });
+    }
+
+    // 2. Try direct path to tokenizer.json
+    if Path::new(model_id).exists() && model_id.ends_with("tokenizer.json") {
+        return Tokenizer::from_file(model_id)
+            .map_err(|e| BenchError::Tokenizer(format!("Failed to load tokenizer: {e}")));
+    }
+
+    // 3. If model_id is a local directory, don't try HF Hub — let tiktoken fallback handle it
+    if Path::new(model_id).is_dir() {
+        return Err(BenchError::Tokenizer(format!(
+            "No tokenizer.json in local directory '{model_id}'"
+        )));
+    }
+
+    // 4. Download from HuggingFace Hub (hf-hub handles auth via HF_TOKEN / cached token)
+    let repo = crate::hub::HubRepo::model(model_id.to_string());
+    let tokenizer_path = repo
+        .get("tokenizer.json")
+        .map_err(|e| BenchError::Tokenizer(format!("No tokenizer.json for '{model_id}': {e}")))?;
+
+    Tokenizer::from_file(&tokenizer_path)
+        .map_err(|e| BenchError::Tokenizer(format!("Failed to load downloaded tokenizer: {e}")))
+}
@@ -173,12 +173,7 @@ impl ResolvedPlaceholder {
        let marker_token_id = context.tokenizer().token_to_id(&token).ok_or_else(|| {
            multimodal!("placeholder token `{token}` is not in the tokenizer vocabulary")
        })?;
-        let embed_token_id = raw.placeholder_token_id_for(&metadata, modality)?;
-        let embed_token_id = u32::try_from(embed_token_id).map_err(|_| {
-            multimodal!(
-                "placeholder token id `{embed_token_id}` for `{modality}` must be non-negative"
-            )
-        })?;
+        let embed_token_id = raw.placeholder_token_id_for(&metadata, modality)? as u32;

        Ok(Self {
            token,
@@ -38,7 +38,7 @@ pub(super) fn build_batched_items(
            let keep_on_cpu = spec.keep_on_cpu_keys.contains(key);
            let (value, field) = match spec.field_layout_for(key) {
                Some(FieldLayout::Batched) => (
-                    tensor.batched_value_at(index)?,
+                    tensor.batched_wire_value_at(index)?,
                    MmField::Batched(MmBatchedField { keep_on_cpu }),
                ),
                Some(FieldLayout::Flat { sizes_key }) => {
@@ -47,7 +47,7 @@ pub(super) fn build_batched_items(
                    })?;
                    let (start, end) = tensor::flat_range_for_index(sizes, sizes_key, index)?;
                    (
-                        tensor.flat_value_range(start, end)?,
+                        tensor.flat_wire_value_range(start, end)?,
                        MmField::Flat(MmFlatField {
                            slices: vec![MmSlice::Slice(SliceSpec {
                                start: Some(0),
@@ -60,7 +60,7 @@ pub(super) fn build_batched_items(
                    )
                }
                None => (
-                    tensor.clone(),
+                    tensor.try_into()?,
                    MmField::Shared(MmSharedField {
                        batch_size: len,
                        keep_on_cpu,
@@ -71,7 +71,7 @@ pub(super) fn build_batched_items(
            data.insert(
                key.clone(),
                MmFieldElem {
-                    data: Some(value.try_into()?),
+                    data: Some(value),
                    field,
                },
            );
--- a/Show More
+++ b/Show More