minor refactor

Signed-off-by: Bugen Zhao <i@bugenzhao.com>
do not take &&, better error categorization
2026-07-24 09:43:05 +00:00 · 2026-07-24 09:26:14 +00:00 · 2026-07-24 09:22:01 +00:00 · 2026-07-24 09:10:46 +00:00 · 2026-07-23 10:16:32 +00:00 · 2026-07-23 07:17:24 +00:00
774 changed files with 42333 additions and 9193 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:
@@ -17,7 +17,7 @@ DEFAULT_REPO_SLUG="vllm-project/vllm"
 DEFAULT_CI_HCL_SOURCE="docker/ci-rocm.hcl"
 DEFAULT_CI_BASE_CONTENT_FILES="requirements/common.txt requirements/rocm.txt requirements/test/rocm.txt docker/Dockerfile.rocm_base docker/ci-rocm.hcl docker/docker-bake-rocm.hcl tools/install_torchcodec_rocm.sh tools/install_protoc.sh rust-toolchain.toml tests/vllm_test_utils .buildkite/scripts/ci-bake-rocm.sh .buildkite/scripts/rocm/build-ci-base.sh"
 DEFAULT_CI_BASE_DOCKERFILE="docker/Dockerfile.rocm"
-DEFAULT_CI_BASE_DOCKERFILE_STAGES="base rust_toolchain_input_0 rust_toolchain_input_1 rust-toolchain-input rust-toolchain build_rixl build_rocshmem build_deepep mori_base ci_base"
+DEFAULT_CI_BASE_DOCKERFILE_STAGES="base rust_toolchain_input_0 rust_toolchain_input_1 rust-toolchain-input rust-toolchain build_nixl build_rocshmem build_deepep mori_base ci_base"
 DEFAULT_CI_BASE_METADATA_VERSION="1"
 IMAGE_EXISTED_BEFORE_BUILD=0

@@ -1159,8 +1159,8 @@ ci_base_metadata_pairs() {
    metadata_pair "vllm.rocm.nic_backend" "$(resolve_dockerfile_arg_value "${dockerfile}" "NIC_BACKEND")"
    metadata_pair "vllm.rocm.ainic_version" "$(resolve_dockerfile_arg_value "${dockerfile}" "AINIC_VERSION")"
    metadata_pair "vllm.rocm.ubuntu_codename" "$(resolve_dockerfile_arg_value "${dockerfile}" "UBUNTU_CODENAME")"
-    metadata_pair "vllm.rocm.rixl_repo" "$(resolve_dockerfile_arg_value "${dockerfile}" "RIXL_REPO")"
-    metadata_pair "vllm.rocm.rixl_commit" "${RIXL_BRANCH:-$(resolve_dockerfile_arg_value "${dockerfile}" "RIXL_BRANCH")}"
+    metadata_pair "vllm.rocm.nixl_repo" "$(resolve_dockerfile_arg_value "${dockerfile}" "NIXL_REPO")"
+    metadata_pair "vllm.rocm.nixl_commit" "${NIXL_BRANCH:-$(resolve_dockerfile_arg_value "${dockerfile}" "NIXL_BRANCH")}"
    metadata_pair "vllm.rocm.ucx_repo" "$(resolve_dockerfile_arg_value "${dockerfile}" "UCX_REPO")"
    metadata_pair "vllm.rocm.ucx_commit" "${UCX_BRANCH:-$(resolve_dockerfile_arg_value "${dockerfile}" "UCX_BRANCH")}"
    metadata_pair "vllm.rocm.rocshmem_repo" "$(resolve_dockerfile_arg_value "${dockerfile}" "ROCSHMEM_REPO")"
@@ -1169,7 +1169,7 @@ ci_base_metadata_pairs() {
    metadata_pair "vllm.rocm.deepep_commit" "${DEEPEP_BRANCH:-$(resolve_dockerfile_arg_value "${dockerfile}" "DEEPEP_BRANCH")}"
    metadata_pair "vllm.rocm.deepep_nic" "$(resolve_dockerfile_arg_value "${dockerfile}" "DEEPEP_NIC")"
    metadata_pair "vllm.rocm.deepep_rocm_arch" "$(resolve_dockerfile_arg_value "${dockerfile}" "DEEPEP_ROCM_ARCH")"
-    metadata_pair "vllm.rocm.rixl_cache_key" "${RIXL_CACHE_KEY:-}"
+    metadata_pair "vllm.rocm.nixl_cache_key" "${NIXL_CACHE_KEY:-}"
    metadata_pair "vllm.rocm.rocshmem_cache_key" "${ROCSHMEM_CACHE_KEY:-}"
    metadata_pair "vllm.rocm.deepep_cache_key" "${DEEPEP_CACHE_KEY:-}"

@@ -1686,7 +1686,7 @@ extract_dependency_pins() {
        return 0
    fi

-    for var in RIXL_BRANCH UCX_BRANCH ROCSHMEM_BRANCH DEEPEP_BRANCH; do
+    for var in NIXL_BRANCH UCX_BRANCH ROCSHMEM_BRANCH DEEPEP_BRANCH; do
        if [[ -n "${!var:-}" ]]; then
            echo "Using provided ${var}: ${!var}"
            continue
@@ -1706,30 +1706,30 @@ extract_dependency_pins() {
 compute_dependency_cache_keys() {
    local bake_dir=""
    local dockerfile_rocm=""
-    local rixl_branch=""
+    local nixl_branch=""
    local ucx_branch=""
    local rocshmem_branch=""
    local deepep_branch=""
-    local rixl_material=""
+    local nixl_material=""
    local rocshmem_material=""
    local deepep_material=""

    bake_dir=$(dirname "${VLLM_BAKE_FILE}")
    dockerfile_rocm="${bake_dir}/Dockerfile.rocm"
-    rixl_branch=$(resolve_dockerfile_arg_value "${dockerfile_rocm}" "RIXL_BRANCH")
+    nixl_branch=$(resolve_dockerfile_arg_value "${dockerfile_rocm}" "NIXL_BRANCH")
    ucx_branch=$(resolve_dockerfile_arg_value "${dockerfile_rocm}" "UCX_BRANCH")
    rocshmem_branch=$(resolve_dockerfile_arg_value "${dockerfile_rocm}" "ROCSHMEM_BRANCH")
    deepep_branch=$(resolve_dockerfile_arg_value "${dockerfile_rocm}" "DEEPEP_BRANCH")

-    if [[ -n "${rixl_branch}" && -n "${ucx_branch}" ]]; then
-        rixl_material=$(compose_stage_cache_material "${dockerfile_rocm}" "base build_rixl")
-        RIXL_CACHE_KEY=$(
+    if [[ -n "${nixl_branch}" && -n "${ucx_branch}" ]]; then
+        nixl_material=$(compose_stage_cache_material "${dockerfile_rocm}" "base build_nixl")
+        NIXL_CACHE_KEY=$(
            compose_dependency_cache_key \
-                "${rixl_branch}-ucx-${ucx_branch}" \
-                "${rixl_material}"
+                "${nixl_branch}-ucx-${ucx_branch}" \
+                "${nixl_material}"
        )
-        export RIXL_CACHE_KEY
-        echo "RIXL dependency cache key: ${RIXL_CACHE_KEY}"
+        export NIXL_CACHE_KEY
+        echo "NIXL dependency cache key: ${NIXL_CACHE_KEY}"
    fi

    if [[ -n "${rocshmem_branch}" ]]; then
@@ -1780,11 +1780,11 @@ dependency_cache_ref_for_target() {
    local cache_repo="${DOCKERHUB_CACHE_REPO:-rocm/vllm-ci-cache}"

    case "${target}" in
-        rixl-rocm-ci)
-            if [[ -n "${RIXL_CACHE_KEY:-}" ]]; then
-                printf '%s\n' "${cache_repo}:rixl-rocm-${RIXL_CACHE_KEY}"
-            elif [[ -n "${RIXL_BRANCH:-}" ]]; then
-                printf '%s\n' "${cache_repo}:rixl-rocm-${RIXL_BRANCH}-ucx-${UCX_BRANCH:-}"
+        nixl-rocm-ci)
+            if [[ -n "${NIXL_CACHE_KEY:-}" ]]; then
+                printf '%s\n' "${cache_repo}:nixl-rocm-${NIXL_CACHE_KEY}"
+            elif [[ -n "${NIXL_BRANCH:-}" ]]; then
+                printf '%s\n' "${cache_repo}:nixl-rocm-${NIXL_BRANCH}-ucx-${UCX_BRANCH:-}"
            fi
            ;;
        rocshmem-rocm-ci)
@@ -1815,7 +1815,7 @@ add_dependency_cache_target() {

 resolve_ci_base_dependency_targets() {
    local mode="${ROCM_DEP_CACHE_EXPORT_MODE:-missing}"
-    local rixl_ref=""
+    local nixl_ref=""
    local rocshmem_ref=""
    local deepep_ref=""

@@ -1824,7 +1824,7 @@ resolve_ci_base_dependency_targets() {
    case "${mode}" in
        always)
            echo "ROCM_DEP_CACHE_EXPORT_MODE=always; exporting all dependency caches serially"
-            for target in rixl-rocm-ci rocshmem-rocm-ci deepep-rocm-ci; do
+            for target in nixl-rocm-ci rocshmem-rocm-ci deepep-rocm-ci; do
                if [[ -n "$(dependency_cache_ref_for_target "${target}")" ]]; then
                    add_dependency_cache_target "${target}"
                fi
@@ -1844,13 +1844,13 @@ resolve_ci_base_dependency_targets() {
            ;;
    esac

-    if [[ "${mode}" != "always" && -n "${RIXL_CACHE_KEY:-}" ]]; then
-        rixl_ref=$(dependency_cache_ref_for_target "rixl-rocm-ci")
-        if dependency_cache_ref_exists "${rixl_ref}"; then
-            echo "RIXL dependency cache exists: ${rixl_ref}"
+    if [[ "${mode}" != "always" && -n "${NIXL_CACHE_KEY:-}" ]]; then
+        nixl_ref=$(dependency_cache_ref_for_target "nixl-rocm-ci")
+        if dependency_cache_ref_exists "${nixl_ref}"; then
+            echo "NIXL dependency cache exists: ${nixl_ref}"
        else
-            echo "RIXL dependency cache missing; will seed: ${rixl_ref}"
-            add_dependency_cache_target "rixl-rocm-ci"
+            echo "NIXL dependency cache missing; will seed: ${nixl_ref}"
+            add_dependency_cache_target "nixl-rocm-ci"
        fi
    fi

@@ -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
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Basic Correctness
  key: basic-correctness
-  timeout_in_minutes: 45
+  timeout_in_minutes: 68
  device: h200_18gb
  source_file_dependencies:
  - vllm/
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Benchmarks CLI Test
  key: benchmarks-cli-test
-  timeout_in_minutes: 30
+  timeout_in_minutes: 45
  device: h200_18gb
  source_file_dependencies:
  - vllm/
@@ -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:
@@ -25,7 +26,10 @@ steps:
  - vllm/v1/cudagraph_dispatcher.py
  - vllm/config/compilation.py
  - vllm/compilation
+  - vllm/v1/worker/encoder_cudagraph.py
+  - vllm/v1/worker/encoder_cudagraph_defs.py
  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
+    - pytest -v -s v1/cudagraph/test_encoder_cudagraph.py
@@ -51,7 +51,7 @@ steps:

 - label: e2e Scheduling (1 GPU)
  key: e2e-scheduling-1-gpu
-  timeout_in_minutes: 35
+  timeout_in_minutes: 53
  device: h200_18gb
  source_file_dependencies:
    - vllm/v1/
@@ -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"
@@ -37,7 +39,7 @@ steps:
 - label: Entrypoints Integration (API Server)
  key: entrypoints-integration-api-server
  device: h200_35gb
-  timeout_in_minutes: 50
+  timeout_in_minutes: 75
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
  - vllm/
@@ -55,8 +57,9 @@ 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
+  timeout_in_minutes: 68
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
  - vllm/
@@ -73,8 +76,9 @@ 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
+  timeout_in_minutes: 83
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
  - vllm/
@@ -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,8 +154,9 @@ steps:
  - pytest -v -s entrypoints/multimodal

 - label: Entrypoints Integration (Pooling)
+  device: h200_35gb
  key: entrypoints-integration-pooling
-  timeout_in_minutes: 50
+  timeout_in_minutes: 75
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
  - vllm/
@@ -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
@@ -31,7 +31,7 @@ steps:

 - label: V1 Sample + Logits
  key: v1-sample-logits
-  timeout_in_minutes: 45
+  timeout_in_minutes: 83
  device: h200_18gb
  source_file_dependencies:
    - vllm/config/
@@ -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/
@@ -89,6 +90,7 @@ steps:
    - tests/v1/kv_offload
    - tests/v1/simple_kv_offload
    - tests/v1/worker
+    - tests/v1/streaming_input
    - tests/v1/kv_connector/unit
    - tests/v1/ec_connector/unit
    - tests/v1/metrics
@@ -102,6 +104,7 @@ steps:
    - pytest -v -s v1/kv_offload
    - pytest -v -s v1/simple_kv_offload
    - pytest -v -s v1/worker
+    - pytest -v -s v1/streaming_input
    - pytest -v -s -m 'not cpu_test' v1/kv_connector/unit
    - pytest -v -s -m 'not cpu_test' v1/ec_connector/unit
    - pytest -v -s -m 'not cpu_test' v1/metrics
@@ -142,6 +145,7 @@ steps:
    - pytest -v -s -m 'cpu_test' v1/core
    - pytest -v -s v1/structured_output
    - pytest -v -s v1/test_serial_utils.py
+    - pytest -v -s v1/cudagraph/test_cudagraph_manager.py
    - pytest -v -s -m 'cpu_test' v1/kv_connector/unit
    - pytest -v -s -m 'cpu_test' v1/metrics

@@ -3,13 +3,16 @@ depends_on:
  - image-build
 steps:
 - label: Model Executor
+  device: h200_35gb
  key: model-executor
-  timeout_in_minutes: 45
+  timeout_in_minutes: 60
  source_file_dependencies:
  - vllm/engine/arg_utils.py
  - vllm/config/model.py
  - vllm/model_executor
+  - vllm/model_executor/warmup
  - tests/model_executor
+  - tests/model_executor/test_jit_warmup.py
  - tests/entrypoints/openai/completion/test_tensorizer_entrypoint.py
  commands:
    - apt-get update && apt-get install -y curl libsodium23
@@ -33,7 +36,9 @@ steps:
      - vllm/engine/arg_utils.py
      - vllm/config/model.py
      - vllm/model_executor
+      - vllm/model_executor/warmup
      - tests/model_executor
+      - tests/model_executor/test_jit_warmup.py
      - tests/entrypoints/openai/completion/test_tensorizer_entrypoint.py
      - vllm/_aiter_ops.py
      - vllm/platforms/rocm.py
@@ -46,6 +46,19 @@ steps:
      depends_on:
      - image-build-amd

+- label: Inkling Unit Tests (B200)
+  key: inkling-unit-tests-b200
+  timeout_in_minutes: 40
+  device: b200-k8s
+  source_file_dependencies:
+  - vllm/models/inkling/
+  - vllm/cute_utils/
+  - cmake/external_projects/tml_fa4.cmake
+  - tests/models/inkling/
+  commands:
+    # FA4 kernel tests require SM100; the suite skips them elsewhere.
+    - pytest -v -s models/inkling
+
 - label: Basic Models Test (Other CPU) # 5min
  key: basic-models-test-other-cpu
  depends_on:
@@ -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
@@ -122,7 +137,7 @@ steps:

 - label: Language Models Test (MTEB)
  key: language-models-test-mteb
-  timeout_in_minutes: 45
+  timeout_in_minutes: 68
  device: h200_18gb
  optional: true
  source_file_dependencies:
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: "Multi-Modal Models (Standard) 1: qwen2"
  key: multi-modal-models-standard-1-qwen2
-  timeout_in_minutes: 45
+  timeout_in_minutes: 68
  device: h200_18gb
  source_file_dependencies:
  - vllm/
@@ -20,7 +20,7 @@ steps:

 - label: "Multi-Modal Models (Standard) 2: qwen3 + gemma"
  key: multi-modal-models-standard-2-qwen3-gemma
-  timeout_in_minutes: 50
+  timeout_in_minutes: 75
  device: h200_18gb
  source_file_dependencies:
  - vllm/
@@ -54,7 +54,7 @@ steps:
 - label: "Multi-Modal Models (Standard) 4: other + whisper"
  device: h200_35gb
  key: multi-modal-models-standard-4-other-whisper
-  timeout_in_minutes: 50
+  timeout_in_minutes: 75
  source_file_dependencies:
  - vllm/
  - tests/models/multimodal
@@ -85,7 +85,7 @@ steps:

 - label: Multi-Modal Processor # 44min
  key: multi-modal-processor
-  timeout_in_minutes: 65
+  timeout_in_minutes: 98
  device: h200_18gb
  source_file_dependencies:
  - vllm/
@@ -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:
@@ -5,7 +5,7 @@ steps:
 - label: PyTorch Compilation Unit Tests
  device: h200_35gb
  key: pytorch-compilation-unit-tests
-  timeout_in_minutes: 90
+  timeout_in_minutes: 110
  source_file_dependencies:
    - vllm/__init__.py
    - vllm/_aiter_ops.py
@@ -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:
@@ -170,3 +170,19 @@ steps:
    - tests/v1/e2e/spec_decode/
  commands:
    - pytest -v -s v1/e2e/spec_decode -k "qwen3_5-hybrid"
+
+- label: Spec Decode DeepSeek MTP Parallel Load (B200)
+  key: spec-decode-deepseek-mtp-parallel-load-b200
+  timeout_in_minutes: 30
+  device: b200-k8s
+  optional: true
+  num_devices: 2
+  source_file_dependencies:
+    - vllm/v1/spec_decode/llm_base_proposer.py
+    - vllm/v1/spec_decode/eagle.py
+    - vllm/v1/worker/gpu/spec_decode/eagle/
+    - vllm/model_executor/models/deepseek_mtp.py
+    - vllm/model_executor/models/deepseek_v2.py
+    - tests/v1/e2e/spec_decode/test_mtp_parallel_load.py
+  commands:
+    - pytest -v -s v1/e2e/spec_decode/test_mtp_parallel_load.py
@@ -47,6 +47,7 @@

 # Rust Frontend
 /rust/ @BugenZhao @njhill
+/rust/src/bench @esmeetu
 /build_rust.sh @BugenZhao @njhill
 /rust-toolchain.toml @BugenZhao @njhill
 /.buildkite/test_areas/rust* @BugenZhao @njhill
@@ -181,6 +181,18 @@ pull_request_rules:
      add:
        - performance

+- name: label-quantization
+  description: Automatically apply quantization label
+  conditions:
+    - label != stale
+    - or:
+      - files~=^vllm/model_executor/layers/quantization/
+      - title~=(?i)quant
+  actions:
+    label:
+      add:
+        - quantization
+
 - name: label-qwen
  description: Automatically apply qwen label
  conditions:
@@ -130,6 +130,47 @@ jobs:
                  },
                ],
              },
+              quantization: {
+                keywords: [
+                  {
+                    term: "quantization",
+                    searchIn: "both"
+                  },
+                  {
+                    term: "quantized",
+                    searchIn: "both"
+                  },
+                ],
+              },
+              "intel-gpu": {
+                // Keyword search - matches whole words only (with word boundaries)
+                keywords: [
+                  {
+                    term: "B50",
+                    searchIn: "both"
+                  },
+                  {
+                    term: "B60",
+                    searchIn: "both"
+                  },
+                  {
+                    term: "B70",
+                    searchIn: "both"
+                  },
+                  {
+                    term: "intel gpu",
+                    searchIn: "both"
+                  },
+                  {
+                    term: "Arc GPU",
+                    searchIn: "both"
+                  },
+                  {
+                    term: "BMG",
+                    searchIn: "both"
+                  },
+                ],
+              },
              // Add more label configurations here as needed
              // example: {
              //   keywords: [...],
@@ -491,4 +532,4 @@ jobs:
              issue_number: context.issue.number,
              body: message,
            });
-            core.notice(`Requested missing ROCm info from @${author}: ${missing.map(m => m.name).join(', ')}`);
+            core.notice(`Requested missing ROCm info from @${author}: ${missing.map(m => m.name).join(', ')}`);
@@ -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:
@@ -48,7 +48,7 @@ vLLM is flexible and easy to use with:
 - Tool calling and reasoning parsers
 - OpenAI-compatible API server, plus Anthropic Messages API and gRPC support
 - Efficient multi-LoRA support for dense and MoE layers
- Support for NVIDIA GPUs, AMD GPUs, and x86/ARM/PowerPC CPUs. Additionally, diverse hardware plugins such as Google TPUs, Intel Gaudi, IBM Spyre, Huawei Ascend, Rebellions NPU, Apple Silicon, MetaX GPU, and more.
+- Support for NVIDIA GPUs, AMD GPUs, Intel GPUs, and x86/ARM/PowerPC CPUs. Additionally, diverse hardware plugins such as Google TPUs, Intel Gaudi, IBM Spyre, Huawei Ascend, Rebellions NPU, Apple Silicon, MetaX GPU, and more.

 vLLM seamlessly supports 200+ model architectures on Hugging Face, including:

@@ -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
@@ -17,7 +17,7 @@ else()
  FetchContent_Declare(
    fmha_sm100
    GIT_REPOSITORY https://github.com/vllm-project/MSA.git
-    GIT_TAG 2e63ec37a0fc29bc20f39cd1a52e0f5affc33a73
+    GIT_TAG 890aaa1a37a598ad17ccff0827fea21540d381fa
    GIT_PROGRESS TRUE
    CONFIGURE_COMMAND ""
    BUILD_COMMAND ""
@@ -22,7 +22,7 @@ if(QUTLASS_SRC_DIR)
  set(qutlass_BINARY_DIR "${CMAKE_BINARY_DIR}/qutlass-binary-dir-unused")
 else()
  set(_QUTLASS_UPSTREAM_REPO "https://github.com/IST-DASLab/qutlass.git")
-  set(_QUTLASS_UPSTREAM_TAG "830d2c4537c7396e14a02a46fbddd18b5d107c65")
+  set(_QUTLASS_UPSTREAM_TAG "e74319e3405ce6d71965732880f5dc1f52371f64")

  set(_qutlass_fc_root "${FETCHCONTENT_BASE_DIR}")
  if(NOT _qutlass_fc_root)
@@ -125,8 +125,6 @@ if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.8 AND QUTLASS_ARCHS)
    CUDA_ARCHS "${QUTLASS_ARCHS}"
  )

-  # QuTLASS uses legacy ATen headers and cannot be built with TORCH_TARGET_VERSION.
-  # Keep it as its own extension (registers torch.ops._qutlass_C).
  define_extension_target(
    _qutlass_C
    DESTINATION vllm
@@ -139,9 +137,11 @@ if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.8 AND QUTLASS_ARCHS)
    WITH_SOABI)

  target_compile_definitions(_qutlass_C PRIVATE
-    QUTLASS_DISABLE_PYBIND=1
+    QUTLASS_MINIMAL_BUILD=1
    TARGET_CUDA_ARCH=${QUTLASS_TARGET_CC}
-    CUTLASS_ENABLE_DIRECT_CUDA_DRIVER_CALL=1)
+    CUTLASS_ENABLE_DIRECT_CUDA_DRIVER_CALL=1
+    TORCH_TARGET_VERSION=0x020B000000000000ULL
+    USE_CUDA)

  set_property(SOURCE ${QUTLASS_SOURCES} APPEND PROPERTY COMPILE_OPTIONS
    $<$<COMPILE_LANGUAGE:CUDA>:--expt-relaxed-constexpr --use_fast_math -O3>
@@ -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 "")
@@ -39,7 +39,7 @@ else()
  FetchContent_Declare(
          vllm-flash-attn
          GIT_REPOSITORY https://github.com/vllm-project/flash-attention.git
-          GIT_TAG caaa4eb59845388a20b1f435ecaafb4bd9517ad8
+          GIT_TAG ed4b7342bc8f0489dd9b649d5288867e35fc6a32
          GIT_PROGRESS TRUE
          # Don't share the vllm-flash-attn build between build types
          BINARY_DIR ${CMAKE_BINARY_DIR}/vllm-flash-attn
@@ -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(
@@ -102,7 +102,9 @@ class TileGemm82 {
    kv_cache_t* __restrict__ curr_b = b_tile;

    for (int32_t k = 0; k < dynamic_k_size; ++k) {
-      auto [fp32_b_0_reg, fp32_b_1_reg] = load_b_pair_vec(curr_b);
+      auto fp32_b_regs = load_b_pair_vec(curr_b);
+      auto fp32_b_0_reg = fp32_b_regs.first;
+      auto fp32_b_1_reg = fp32_b_regs.second;

      float* __restrict__ curr_m_a = curr_a;
      vec_op::unroll_loop<int32_t, M>([&](int32_t i) {
@@ -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
@@ -1025,6 +1025,9 @@ __global__ void gather_and_maybe_dequant_cache(
    batch_offset += offset;
    int32_t block_table_id = batch_offset / block_size;
    int32_t slot_id = batch_offset % block_size;
+    // seq_starts may push the block index past the end of the batch's block
+    // table row.
+    if (block_table_id >= block_table_stride) continue;
    int32_t block_table_offset = batch_id * block_table_stride + block_table_id;
    int32_t block_id = block_table[block_table_offset];
    int64_t cache_offset =
@@ -1174,7 +1177,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 +1196,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 +1388,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 +1401,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 +1448,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 +1459,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.
@@ -9,14 +9,16 @@ void topk_softmax(torch::stable::Tensor& topk_weights,
                  torch::stable::Tensor& topk_indices,
                  torch::stable::Tensor& token_expert_indices,
                  torch::stable::Tensor& gating_output, bool renormalize,
-                  std::optional<torch::stable::Tensor> bias);
+                  std::optional<torch::stable::Tensor> bias,
+                  std::optional<torch::stable::Tensor> is_padding);

 void topk_sigmoid(torch::stable::Tensor& topk_weights,
                  torch::stable::Tensor& topk_indices,
                  torch::stable::Tensor& token_expert_indices,
                  torch::stable::Tensor& gating_output, bool renormalize,
                  std::optional<torch::stable::Tensor> bias,
-                  double routed_scaling_factor);
+                  double routed_scaling_factor,
+                  std::optional<torch::stable::Tensor> is_padding);

 void topk_softplus_sqrt(
    torch::stable::Tensor& topk_weights, torch::stable::Tensor& topk_indices,
@@ -25,7 +27,8 @@ void topk_softplus_sqrt(
    double routed_scaling_factor,
    const std::optional<torch::stable::Tensor>& correction_bias,
    const std::optional<torch::stable::Tensor>& input_ids,
-    const std::optional<torch::stable::Tensor>& tid2eid);
+    const std::optional<torch::stable::Tensor>& tid2eid,
+    const std::optional<torch::stable::Tensor>& is_padding);

 void moe_sum(torch::stable::Tensor& input, torch::stable::Tensor& output,
             std::optional<torch::stable::Tensor> topk_ids,
@@ -174,7 +174,8 @@ __launch_bounds__(TPB) __global__ void moeTopK(
    const int end_expert,
    const bool renormalize,
    const float* bias,
-    const double routed_scaling_factor)
+    const double routed_scaling_factor,
+    const bool* is_padding)
 {

    using cub_kvp = cub::KeyValuePair<int, float>;
@@ -228,12 +229,14 @@ __launch_bounds__(TPB) __global__ void moeTopK(
            const int expert = result_kvp.key;
            const bool node_uses_expert = expert >= start_expert && expert < end_expert;
            const bool should_process_row = row_is_active && node_uses_expert;
+            const bool is_pad_row = is_padding != nullptr && is_padding[block_row];

            const int idx = k * block_row + k_idx;
            // Return the unbiased scores for output weights
            output[idx] = inputs_after_softmax[thread_read_offset + expert];
-            indices[idx] = should_process_row ? (expert - start_expert) : num_experts;
-            assert(indices[idx] >= 0);
+            indices[idx] = is_pad_row ? static_cast<IndType>(-1)
+                                       : (should_process_row ? (expert - start_expert) : num_experts);
+            assert(is_pad_row || indices[idx] >= 0);
            source_rows[idx] = k_idx * num_rows + block_row;
            if (renormalize) {
                selected_sum += inputs_after_softmax[thread_read_offset + expert];
@@ -277,7 +280,7 @@ template <int VPT, int NUM_EXPERTS, int WARPS_PER_CTA, int BYTES_PER_LDG, int WA
 __launch_bounds__(WARPS_PER_CTA* WARP_SIZE_PARAM) __global__
    void topkGating(const InputType* input, const bool* finished, float* output, const int num_rows, IndType* indices,
        int* source_rows, const int k, const int start_expert, const int end_expert, const bool renormalize,
-        const float* bias, const double routed_scaling_factor)
+        const float* bias, const double routed_scaling_factor, const bool* is_padding)
 {
    static_assert(std::is_same_v<InputType, float> || std::is_same_v<InputType, __nv_bfloat16> ||
                      std::is_same_v<InputType, __half>,
@@ -545,12 +548,14 @@ __launch_bounds__(WARPS_PER_CTA* WARP_SIZE_PARAM) __global__
            // Add a guard to ignore experts not included by this node
            const bool node_uses_expert = expert >= start_expert && expert < end_expert;
            const bool should_process_row = row_is_active && node_uses_expert;
+            const bool is_pad_row = is_padding != nullptr && is_padding[thread_row];

            // The lead thread from each sub-group will write out the final results to global memory. (This will be a
            // single) thread per row of the input/output matrices.
            const int idx = k * thread_row + k_idx;
            output[idx] = max_val;
-            indices[idx] = should_process_row ? (expert - start_expert) : NUM_EXPERTS;
+            indices[idx] = is_pad_row ? static_cast<IndType>(-1)
+                                       : (should_process_row ? (expert - start_expert) : NUM_EXPERTS);
            source_rows[idx] = k_idx * num_rows + thread_row;
            if (renormalize) {
                selected_sum += max_val;
@@ -605,7 +610,7 @@ struct TopkConstants
 template <int EXPERTS, int WARPS_PER_TB, int WARP_SIZE_PARAM, int MAX_BYTES_PER_LDG, typename IndType, typename InputType, ScoringFunc SF>
 void topkGatingLauncherHelper(const InputType* input, const bool* finished, float* output, IndType* indices,
    int* source_row, const int num_rows, const int k, const int start_expert, const int end_expert, const bool renormalize,
-    const float* bias, const double routed_scaling_factor, cudaStream_t stream)
+    const float* bias, const double routed_scaling_factor, cudaStream_t stream, const bool* is_padding)
 {
    static constexpr int BYTES_PER_LDG = MIN(MAX_BYTES_PER_LDG, sizeof(InputType) * EXPERTS);
    using Constants = detail::TopkConstants<EXPERTS, BYTES_PER_LDG, WARP_SIZE_PARAM, InputType>;
@@ -616,7 +621,7 @@ void topkGatingLauncherHelper(const InputType* input, const bool* finished, floa

    dim3 block_dim(WARP_SIZE_PARAM, WARPS_PER_TB);
    topkGating<VPT, EXPERTS, WARPS_PER_TB, BYTES_PER_LDG, WARP_SIZE_PARAM, IndType, InputType, SF><<<num_blocks, block_dim, 0, stream>>>(
-        input, finished, output, num_rows, indices, source_row, k, start_expert, end_expert, renormalize, bias, routed_scaling_factor);
+        input, finished, output, num_rows, indices, source_row, k, start_expert, end_expert, renormalize, bias, routed_scaling_factor, is_padding);
 }

 #ifndef USE_ROCM
@@ -627,7 +632,7 @@ void topkGatingLauncherHelper(const InputType* input, const bool* finished, floa
                             IndType, InputType, SF>(                         \
        gating_output, nullptr, topk_weights, topk_indices,                   \
        token_expert_indices, num_tokens, topk, 0, num_experts, renormalize,  \
-        bias, routed_scaling_factor, stream);
+        bias, routed_scaling_factor, stream, is_padding);
 #else
  #define LAUNCH_TOPK(NUM_EXPERTS, WARPS_PER_TB, MAX_BYTES)                    \
    if (WARP_SIZE == 64) {                                                     \
@@ -635,13 +640,13 @@ void topkGatingLauncherHelper(const InputType* input, const bool* finished, floa
                               IndType, InputType, SF>(                        \
          gating_output, nullptr, topk_weights, topk_indices,                  \
          token_expert_indices, num_tokens, topk, 0, num_experts, renormalize, \
-          bias, routed_scaling_factor, stream);                                                       \
+          bias, routed_scaling_factor, stream, is_padding);                                            \
    } else if (WARP_SIZE == 32) {                                              \
      topkGatingLauncherHelper<NUM_EXPERTS, WARPS_PER_TB, 32, MAX_BYTES,       \
                               IndType, InputType, SF>(                        \
          gating_output, nullptr, topk_weights, topk_indices,                  \
          token_expert_indices, num_tokens, topk, 0, num_experts, renormalize, \
-          bias, routed_scaling_factor, stream);                                                       \
+          bias, routed_scaling_factor, stream, is_padding);                                            \
    } else {                                                                   \
      assert(false &&                                                          \
             "Unsupported warp size. Only 32 and 64 are supported for ROCm");  \
@@ -661,7 +666,8 @@ void topkGatingKernelLauncher(
    const bool renormalize,
    const float* bias,
    const double routed_scaling_factor,
-    cudaStream_t stream) {
+    cudaStream_t stream,
+    const bool* is_padding) {
    static constexpr int WARPS_PER_TB = 4;
    static constexpr int BYTES_PER_LDG_POWER_OF_2 = 16;
 #ifndef USE_ROCM
@@ -736,7 +742,7 @@ void topkGatingKernelLauncher(
            }
            moeTopK<TPB><<<num_tokens, TPB, 0, stream>>>(
                workspace, nullptr, topk_weights, topk_indices, token_expert_indices,
-                num_experts, topk, 0, num_experts, renormalize, bias, routed_scaling_factor);
+                num_experts, topk, 0, num_experts, renormalize, bias, routed_scaling_factor, is_padding);
        }
    }
 }
@@ -755,7 +761,8 @@ void dispatch_topk_launch(
    int num_tokens, int num_experts, int topk, bool renormalize,
    std::optional<torch::stable::Tensor> bias,
    double routed_scaling_factor,
-    cudaStream_t stream)
+    cudaStream_t stream,
+    std::optional<torch::stable::Tensor> is_padding)
 {
    const float* bias_ptr = nullptr;
    if (bias.has_value()) {
@@ -769,6 +776,18 @@ void dispatch_topk_launch(
      bias_ptr = bias_tensor.const_data_ptr<float>();
    }

+    const bool* is_padding_ptr = nullptr;
+    if (is_padding.has_value()) {
+      const torch::stable::Tensor& is_padding_tensor = is_padding.value();
+      STD_TORCH_CHECK(is_padding_tensor.scalar_type() == torch::headeronly::ScalarType::Bool,
+                      "is_padding tensor must be bool");
+      STD_TORCH_CHECK(is_padding_tensor.dim() == 1, "is_padding tensor must be 1D");
+      STD_TORCH_CHECK(is_padding_tensor.size(0) == num_tokens,
+                      "is_padding size mismatch, expected: ", num_tokens);
+      STD_TORCH_CHECK(is_padding_tensor.is_contiguous(), "is_padding tensor must be contiguous");
+      is_padding_ptr = is_padding_tensor.const_data_ptr<bool>();
+    }
+
    if (topk_indices.scalar_type() == torch::headeronly::ScalarType::Int) {
        vllm::moe::topkGatingKernelLauncher<int, ComputeType, SF>(
            reinterpret_cast<const ComputeType*>(gating_output.const_data_ptr()),
@@ -777,7 +796,7 @@ void dispatch_topk_launch(
            token_expert_indices.mutable_data_ptr<int>(),
            softmax_workspace.mutable_data_ptr<float>(),
            num_tokens, num_experts, topk, renormalize,
-            bias_ptr, routed_scaling_factor, stream);
+            bias_ptr, routed_scaling_factor, stream, is_padding_ptr);
    } else if (topk_indices.scalar_type() == torch::headeronly::ScalarType::UInt32) {
        vllm::moe::topkGatingKernelLauncher<uint32_t, ComputeType, SF>(
            reinterpret_cast<const ComputeType*>(gating_output.const_data_ptr()),
@@ -786,7 +805,7 @@ void dispatch_topk_launch(
            token_expert_indices.mutable_data_ptr<int>(),
            softmax_workspace.mutable_data_ptr<float>(),
            num_tokens, num_experts, topk, renormalize,
-            bias_ptr, routed_scaling_factor, stream);
+            bias_ptr, routed_scaling_factor, stream, is_padding_ptr);
    } else {
        STD_TORCH_CHECK(topk_indices.scalar_type() == torch::headeronly::ScalarType::Long);
        vllm::moe::topkGatingKernelLauncher<int64_t, ComputeType, SF>(
@@ -796,7 +815,7 @@ void dispatch_topk_launch(
            token_expert_indices.mutable_data_ptr<int>(),
            softmax_workspace.mutable_data_ptr<float>(),
            num_tokens, num_experts, topk, renormalize,
-            bias_ptr, routed_scaling_factor, stream);
+            bias_ptr, routed_scaling_factor, stream, is_padding_ptr);
    }
 }

@@ -806,7 +825,8 @@ void topk_softmax(
    torch::stable::Tensor& token_expert_indices,        // [num_tokens, topk]
    torch::stable::Tensor& gating_output,               // [num_tokens, num_experts]
    bool renormalize,
-    std::optional<torch::stable::Tensor> bias)
+    std::optional<torch::stable::Tensor> bias,
+    std::optional<torch::stable::Tensor> is_padding)
 {
    const int num_experts = gating_output.size(-1);
    const auto num_tokens = gating_output.numel() / num_experts;
@@ -825,15 +845,15 @@ void topk_softmax(
    if (gating_output.scalar_type() == torch::headeronly::ScalarType::Float) {
        dispatch_topk_launch<float, vllm::moe::SCORING_SOFTMAX>(gating_output, topk_weights, topk_indices,
            token_expert_indices, softmax_workspace, num_tokens, num_experts, topk, renormalize,
-            bias, 1.0, stream);
+            bias, 1.0, stream, is_padding);
    } else if (gating_output.scalar_type() == torch::headeronly::ScalarType::Half) {
        dispatch_topk_launch<__half, vllm::moe::SCORING_SOFTMAX>(gating_output, topk_weights, topk_indices,
            token_expert_indices, softmax_workspace, num_tokens, num_experts, topk, renormalize,
-            bias, 1.0, stream);
+            bias, 1.0, stream, is_padding);
    } else if (gating_output.scalar_type() == torch::headeronly::ScalarType::BFloat16) {
        dispatch_topk_launch<__nv_bfloat16, vllm::moe::SCORING_SOFTMAX>(gating_output, topk_weights, topk_indices,
            token_expert_indices, softmax_workspace, num_tokens, num_experts, topk, renormalize,
-            bias, 1.0, stream);
+            bias, 1.0, stream, is_padding);
    } else {
        STD_TORCH_CHECK(false, "Unsupported gating_output data type: ", gating_output.scalar_type());
    }
@@ -846,7 +866,8 @@ void topk_sigmoid(
    torch::stable::Tensor& gating_output,               // [num_tokens, num_experts]
    bool renormalize,
    std::optional<torch::stable::Tensor> bias,
-    double routed_scaling_factor)
+    double routed_scaling_factor,
+    std::optional<torch::stable::Tensor> is_padding)
 {
    const int num_experts = gating_output.size(-1);
    const auto num_tokens = gating_output.numel() / num_experts;
@@ -865,15 +886,15 @@ void topk_sigmoid(
    if (gating_output.scalar_type() == torch::headeronly::ScalarType::Float) {
        dispatch_topk_launch<float, vllm::moe::SCORING_SIGMOID>(gating_output, topk_weights, topk_indices,
            token_expert_indices, workspace, num_tokens, num_experts, topk, renormalize,
-            bias, routed_scaling_factor, stream);
+            bias, routed_scaling_factor, stream, is_padding);
    } else if (gating_output.scalar_type() == torch::headeronly::ScalarType::Half) {
        dispatch_topk_launch<__half, vllm::moe::SCORING_SIGMOID>(gating_output, topk_weights, topk_indices,
            token_expert_indices, workspace, num_tokens, num_experts, topk, renormalize,
-            bias, routed_scaling_factor, stream);
+            bias, routed_scaling_factor, stream, is_padding);
    } else if (gating_output.scalar_type() == torch::headeronly::ScalarType::BFloat16) {
        dispatch_topk_launch<__nv_bfloat16, vllm::moe::SCORING_SIGMOID>(gating_output, topk_weights, topk_indices,
            token_expert_indices, workspace, num_tokens, num_experts, topk, renormalize,
-            bias, routed_scaling_factor, stream);
+            bias, routed_scaling_factor, stream, is_padding);
    } else {
        STD_TORCH_CHECK(false, "Unsupported gating_output data type: ", gating_output.scalar_type());
    }
@@ -71,6 +71,80 @@ __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 bool* is_padding) {
+  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);
+  const bool is_pad_row = is_padding != nullptr && is_padding[warp];
+
+  #if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900)
+  cudaGridDependencySynchronize();
+  #endif
+  int expert = 0;
+  float weight = 0.f;
+  if (lane < 6 && !is_pad_row) {
+    // 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))));
+    if (isnan(weight)) {
+      weight = 0.f;
+    }
+  }
+  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] = !is_pad_row ? static_cast<OutIndType>(expert)
+                                  : static_cast<OutIndType>(-1);
+  }
+}
+
+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,
+                        const bool* is_padding) {
+  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, is_padding);
+}
+#endif
+
 // ====================== TopK softplus_sqrt things
 // ===============================

@@ -99,7 +173,8 @@ __launch_bounds__(WARPS_PER_CTA* WARP_SIZE_PARAM) __global__
        const int num_rows, IndType* indices, int* source_rows, const int k,
        const int start_expert, const int end_expert, const bool renormalize,
        double routed_scaling_factor, const float* correction_bias,
-        const HashIndType* input_ids, const HashIndType* tid2eid) {
+        const HashIndType* input_ids, const HashIndType* tid2eid,
+        const bool* is_padding) {
  static_assert(std::is_same_v<InputType, float> ||
                    std::is_same_v<InputType, __nv_bfloat16> ||
                    std::is_same_v<InputType, __half>,
@@ -164,6 +239,7 @@ __launch_bounds__(WARPS_PER_CTA* WARP_SIZE_PARAM) __global__
    return;
  }
  const bool row_is_active = finished ? !finished[thread_row] : true;
+  const bool is_pad_row = is_padding != nullptr && is_padding[thread_row];

  // We finally start setting up the read pointers for each thread. First, each
  // thread jumps to the start of the row it will read.
@@ -182,9 +258,12 @@ __launch_bounds__(WARPS_PER_CTA* WARP_SIZE_PARAM) __global__
  cudaGridDependencySynchronize();
 #endif

-  // NOTE(zhuhaoran): dispatch different input types loading, BF16/FP16 convert
-  // to float
-  if constexpr (std::is_same_v<InputType, float>) {
+  if (is_pad_row) {
+#pragma unroll
+    for (int ii = 0; ii < VPT; ++ii) {
+      row_chunk[ii] = 0.f;
+    }
+  } else if constexpr (std::is_same_v<InputType, float>) {
    using VecType = AlignedArray<float, ELTS_PER_LDG>;
    VecType* row_chunk_vec_ptr = reinterpret_cast<VecType*>(&row_chunk);
    const VecType* vec_thread_read_ptr =
@@ -248,12 +327,22 @@ __launch_bounds__(WARPS_PER_CTA* WARP_SIZE_PARAM) __global__
  if constexpr (USE_HASH) {
    const int64_t token_id = load_index_as_int64(input_ids, thread_row);
    const int64_t token_expert_offset = token_id * static_cast<int64_t>(k);
+    if (!is_pad_row) {
 #pragma unroll
-    for (int ii = 0; ii < VPT; ++ii) {
-      float val = row_chunk[ii];
-      float val_b = val * beta;
-      val = (val_b > threshold) ? val : (__logf(1.0f + __expf(val_b))) / beta;
-      row_chunk[ii] = sqrtf(val);
+      for (int ii = 0; ii < VPT; ++ii) {
+        float val = row_chunk[ii];
+        float val_b = val * beta;
+        val = (val_b > threshold) ? val : (__logf(1.0f + __expf(val_b))) / beta;
+        val = sqrtf(val);
+
+        // Dummy/padding tokens can result in NaN values, so
+        // clamp them to 0.0. Note: this clamp could likely be removed if
+        // 'is_padding' is made mandatory
+        if (isnan(val)) {
+          val = 0.f;
+        }
+        row_chunk[ii] = val;
+      }
    }
    float selected_sum = 0.f;
 #pragma unroll
@@ -268,7 +357,8 @@ __launch_bounds__(WARPS_PER_CTA* WARP_SIZE_PARAM) __global__
                               group_id * THREADS_PER_ROW * ELTS_PER_LDG +
                               local_id;
        if (expert == expert_idx) {
-          indices[idx] = static_cast<IndType>(expert);
+          indices[idx] = !is_pad_row ? static_cast<IndType>(expert)
+                                     : static_cast<IndType>(-1);
          selected_sum += row_chunk[ii];
          break;
        }
@@ -312,23 +402,31 @@ __launch_bounds__(WARPS_PER_CTA* WARP_SIZE_PARAM) __global__
 #endif
    return;
  } else {
+    if (!is_pad_row) {
 #pragma unroll
-    for (int ii = 0; ii < VPT; ++ii) {
-      float val = row_chunk[ii];
-      float val_b = val * beta;
-      // Compute softplus: log(1 + exp(val)) with numerical stability
-      // When val > threshold, softplus(x) ≈ x to avoid exp overflow
-      val = (val_b > threshold) ? val : (__logf(1.0f + __expf(val_b))) / beta;
-      val = sqrtf(val);
-      if (correction_bias) {
-        const int group_id = ii / ELTS_PER_LDG;
-        const int local_id = ii % ELTS_PER_LDG;
-        const int expert_idx = first_elt_read_by_thread +
-                               group_id * THREADS_PER_ROW * ELTS_PER_LDG +
-                               local_id;
-        val = val + correction_bias[expert_idx];
+      for (int ii = 0; ii < VPT; ++ii) {
+        float val = row_chunk[ii];
+        float val_b = val * beta;
+        // Compute softplus: log(1 + exp(val)) with numerical stability
+        // When val > threshold, softplus(x) ≈ x to avoid exp overflow
+        val = (val_b > threshold) ? val : (__logf(1.0f + __expf(val_b))) / beta;
+        val = sqrtf(val);
+        // Dummy/padding tokens can result in NaN values, so
+        // clamp them to 0.0. Note: this clamp could likely be removed if
+        // 'is_padding' is made mandatory
+        if (isnan(val)) {
+          val = 0.f;
+        }
+        if (correction_bias) {
+          const int group_id = ii / ELTS_PER_LDG;
+          const int local_id = ii % ELTS_PER_LDG;
+          const int expert_idx = first_elt_read_by_thread +
+                                 group_id * THREADS_PER_ROW * ELTS_PER_LDG +
+                                 local_id;
+          val = val + correction_bias[expert_idx];
+        }
+        row_chunk[ii] = val;
      }
-      row_chunk[ii] = val;
    }

    // Original TopK path: find top-k experts by score
@@ -383,18 +481,19 @@ __launch_bounds__(WARPS_PER_CTA* WARP_SIZE_PARAM) __global__
        // Add a guard to ignore experts not included by this node
        const bool node_uses_expert =
            expert >= start_expert && expert < end_expert;
-        const bool should_process_row = row_is_active && node_uses_expert;
+        const bool should_process_row =
+            row_is_active && node_uses_expert && !is_pad_row;

        // The lead thread from each sub-group will write out the final results
        // to global memory. (This will be a single) thread per row of the
        // input/output matrices.
        const int idx = k * thread_row + k_idx;
-        if (correction_bias != nullptr) {
+        if (correction_bias != nullptr && should_process_row) {
          max_val -= correction_bias[expert];
        }
        output[idx] = max_val;
        indices[idx] =
-            should_process_row ? (expert - start_expert) : NUM_EXPERTS;
+            !is_pad_row ? expert - start_expert : static_cast<IndType>(-1);
        source_rows[idx] = k_idx * num_rows + thread_row;
        if (renormalize) {
          selected_sum += max_val;
@@ -477,7 +576,7 @@ void topkGatingSoftplusSqrtLauncherHelper(
    const int start_expert, const int end_expert, const bool renormalize,
    double routed_scaling_factor, const float* correction_bias,
    const bool use_hash, const HashIndType* input_ids,
-    const HashIndType* tid2eid, cudaStream_t stream) {
+    const HashIndType* tid2eid, cudaStream_t stream, const bool* is_padding) {
  static constexpr int BYTES_PER_LDG =
      MIN(MAX_BYTES_PER_LDG, sizeof(InputType) * EXPERTS);
  using Constants =
@@ -506,12 +605,12 @@ void topkGatingSoftplusSqrtLauncherHelper(
    cudaLaunchKernelEx(&config, kernel, input, finished, output, num_rows,
                       indices, source_row, k, start_expert, end_expert,
                       renormalize, routed_scaling_factor, correction_bias,
-                       input_ids, tid2eid);
+                       input_ids, tid2eid, is_padding);
 #else
    kernel<<<num_blocks, block_dim, 0, stream>>>(
        input, finished, output, num_rows, indices, source_row, k, start_expert,
        end_expert, renormalize, routed_scaling_factor, correction_bias,
-        input_ids, tid2eid);
+        input_ids, tid2eid, is_padding);
 #endif
  })
 }
@@ -525,7 +624,7 @@ void topkGatingSoftplusSqrtLauncherHelper(
        gating_output, nullptr, topk_weights, topk_indices,                    \
        token_expert_indices, num_tokens, topk, 0, num_experts, renormalize,   \
        routed_scaling_factor, correction_bias, use_hash, input_ids, tid2eid,  \
-        stream);
+        stream, is_padding);
 #else
  #define LAUNCH_SOFTPLUS_SQRT(NUM_EXPERTS, WARPS_PER_TB, MAX_BYTES)           \
    if (WARP_SIZE == 64) {                                                     \
@@ -534,14 +633,14 @@ void topkGatingSoftplusSqrtLauncherHelper(
          gating_output, nullptr, topk_weights, topk_indices,                  \
          token_expert_indices, num_tokens, topk, 0, num_experts, renormalize, \
          routed_scaling_factor, correction_bias, use_hash, input_ids,         \
-          tid2eid, stream);                                                    \
+          tid2eid, stream, is_padding);                                        \
    } else if (WARP_SIZE == 32) {                                              \
      topkGatingSoftplusSqrtLauncherHelper<NUM_EXPERTS, WARPS_PER_TB, 32,      \
                                           MAX_BYTES>(                         \
          gating_output, nullptr, topk_weights, topk_indices,                  \
          token_expert_indices, num_tokens, topk, 0, num_experts, renormalize, \
          routed_scaling_factor, correction_bias, use_hash, input_ids,         \
-          tid2eid, stream);                                                    \
+          tid2eid, stream, is_padding);                                        \
    } else {                                                                   \
      assert(false &&                                                          \
             "Unsupported warp size. Only 32 and 64 are supported for ROCm");  \
@@ -555,7 +654,18 @@ void topkGatingSoftplusSqrtKernelLauncher(
    const int topk, const bool renormalize, double routed_scaling_factor,
    const float* correction_bias, const bool use_hash,
    const HashIndType* input_ids, const HashIndType* tid2eid,
-    cudaStream_t stream) {
+    cudaStream_t stream, const bool* is_padding) {
+#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, is_padding);
+      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
@@ -650,7 +760,8 @@ void dispatch_topk_softplus_sqrt_launch(
    int num_experts, int topk, bool renormalize, double routed_scaling_factor,
    const std::optional<torch::stable::Tensor>& correction_bias,
    const std::optional<torch::stable::Tensor>& input_ids,
-    const std::optional<torch::stable::Tensor>& tid2eid, cudaStream_t stream) {
+    const std::optional<torch::stable::Tensor>& tid2eid, cudaStream_t stream,
+    const std::optional<torch::stable::Tensor>& is_padding) {
  const float* bias_ptr = nullptr;
  if (correction_bias.has_value()) {
    bias_ptr = correction_bias.value().const_data_ptr<float>();
@@ -659,6 +770,22 @@ void dispatch_topk_softplus_sqrt_launch(
  auto launch = [&](auto* topk_indices_ptr) {
    using OutIndType =
        typename std::remove_pointer<decltype(topk_indices_ptr)>::type;
+
+    const bool* is_padding_ptr = nullptr;
+    if (is_padding.has_value()) {
+      const torch::stable::Tensor& is_padding_tensor = is_padding.value();
+      STD_TORCH_CHECK(is_padding_tensor.scalar_type() ==
+                          torch::headeronly::ScalarType::Bool,
+                      "is_padding tensor must be bool");
+      STD_TORCH_CHECK(is_padding_tensor.dim() == 1,
+                      "is_padding tensor must be 1D");
+      STD_TORCH_CHECK(is_padding_tensor.size(0) == num_tokens,
+                      "is_padding size mismatch, expected: ", num_tokens);
+      STD_TORCH_CHECK(is_padding_tensor.is_contiguous(),
+                      "is_padding tensor must be contiguous");
+      is_padding_ptr = is_padding_tensor.const_data_ptr<bool>();
+    }
+
    if (tid2eid.has_value()) {
      STD_TORCH_CHECK(input_ids.has_value(),
                      "input_ids is required for hash MoE");
@@ -673,7 +800,7 @@ void dispatch_topk_softplus_sqrt_launch(
            topk_indices_ptr, token_expert_indices.mutable_data_ptr<int>(),
            num_tokens, num_experts, topk, renormalize, routed_scaling_factor,
            bias_ptr, true, input_ids.value().const_data_ptr<int64_t>(),
-            tid2eid.value().const_data_ptr<int64_t>(), stream);
+            tid2eid.value().const_data_ptr<int64_t>(), stream, is_padding_ptr);
      } else {
        STD_TORCH_CHECK(tid2eid.value().scalar_type() ==
                        torch::headeronly::ScalarType::Int);
@@ -683,7 +810,7 @@ void dispatch_topk_softplus_sqrt_launch(
            topk_indices_ptr, token_expert_indices.mutable_data_ptr<int>(),
            num_tokens, num_experts, topk, renormalize, routed_scaling_factor,
            bias_ptr, true, input_ids.value().const_data_ptr<int>(),
-            tid2eid.value().const_data_ptr<int>(), stream);
+            tid2eid.value().const_data_ptr<int>(), stream, is_padding_ptr);
      }
    } else {
      vllm::moe::topkGatingSoftplusSqrtKernelLauncher<OutIndType, ComputeType>(
@@ -691,7 +818,7 @@ void dispatch_topk_softplus_sqrt_launch(
          topk_indices_ptr, token_expert_indices.mutable_data_ptr<int>(),
          num_tokens, num_experts, topk, renormalize, routed_scaling_factor,
          bias_ptr, false, static_cast<const OutIndType*>(nullptr),
-          static_cast<const OutIndType*>(nullptr), stream);
+          static_cast<const OutIndType*>(nullptr), stream, is_padding_ptr);
    }
  };

@@ -715,7 +842,8 @@ void topk_softplus_sqrt(
    bool renormalize, double routed_scaling_factor,
    const std::optional<torch::stable::Tensor>& correction_bias,
    const std::optional<torch::stable::Tensor>& input_ids,
-    const std::optional<torch::stable::Tensor>& tid2eid) {
+    const std::optional<torch::stable::Tensor>& tid2eid,
+    const std::optional<torch::stable::Tensor>& is_padding) {
  const int num_experts = gating_output.size(-1);
  const auto num_tokens = gating_output.numel() / num_experts;
  const int topk = topk_weights.size(-1);
@@ -728,21 +856,22 @@ void topk_softplus_sqrt(
    dispatch_topk_softplus_sqrt_launch<float>(
        gating_output.const_data_ptr<float>(), topk_weights, topk_indices,
        token_expert_indices, num_tokens, num_experts, topk, renormalize,
-        routed_scaling_factor, correction_bias, input_ids, tid2eid, stream);
+        routed_scaling_factor, correction_bias, input_ids, tid2eid, stream,
+        is_padding);
  } else if (gating_output.scalar_type() ==
             torch::headeronly::ScalarType::Half) {
    dispatch_topk_softplus_sqrt_launch<__half>(
        reinterpret_cast<const __half*>(gating_output.const_data_ptr()),
        topk_weights, topk_indices, token_expert_indices, num_tokens,
        num_experts, topk, renormalize, routed_scaling_factor, correction_bias,
-        input_ids, tid2eid, stream);
+        input_ids, tid2eid, stream, is_padding);
  } else if (gating_output.scalar_type() ==
             torch::headeronly::ScalarType::BFloat16) {
    dispatch_topk_softplus_sqrt_launch<__nv_bfloat16>(
        reinterpret_cast<const __nv_bfloat16*>(gating_output.const_data_ptr()),
        topk_weights, topk_indices, token_expert_indices, num_tokens,
        num_experts, topk, renormalize, routed_scaling_factor, correction_bias,
-        input_ids, tid2eid, stream);
+        input_ids, tid2eid, stream, is_padding);
  } else {
    STD_TORCH_CHECK(false, "Unsupported gating_output data type: ",
                    gating_output.scalar_type());
@@ -8,19 +8,19 @@ STABLE_TORCH_LIBRARY_FRAGMENT(_moe_C, m) {
  m.def(
      "topk_softmax(Tensor! topk_weights, Tensor! topk_indices, Tensor! "
      "token_expert_indices, Tensor gating_output, bool renormalize, Tensor? "
-      "bias) -> ()");
+      "bias, Tensor? is_padding) -> ()");

  // Apply topk sigmoid to the gating outputs.
  m.def(
      "topk_sigmoid(Tensor! topk_weights, Tensor! topk_indices, Tensor! "
      "token_expert_indices, Tensor gating_output, bool renormalize, "
-      "Tensor? bias, float routed_scaling_factor) -> ()");
+      "Tensor? bias, float routed_scaling_factor, Tensor? is_padding) -> ()");

  m.def(
      "topk_softplus_sqrt(Tensor! topk_weights, Tensor! topk_indices, Tensor! "
      "token_expert_indices, Tensor gating_output, bool renormalize, float "
      "routed_scaling_factor, Tensor? "
-      "bias, Tensor? input_ids, Tensor? tid2eid) -> ()");
+      "bias, Tensor? input_ids, Tensor? tid2eid, Tensor? is_padding) -> ()");

  // Calculate the result of moe by summing up the partial results
  // from all selected experts. topk_ids/expert_map are optional and, when
@@ -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(
@@ -39,11 +39,15 @@ __global__ void marlin_int4_fp8_preprocess_kernel_awq(
    // AWQ zeros: (size_k // group_size, size_n // 8)
    const int32_t* __restrict__ qzeros, int32_t size_n, int32_t size_k,
    int32_t group_size) {
-  int32_t val =
-      qweight[(blockIdx.x * 32 + threadIdx.x) * size_n / 8 + blockIdx.y];
-  int32_t zero =
-      qzeros[(blockIdx.x * 32 + threadIdx.x) / group_size * size_n / 8 +
-             blockIdx.y];
+  // Thread mapping: threadIdx.x -> column dim (coalesced read within a row),
+  // blockIdx.x -> row dim. Adjacent threads read consecutive int32 in the
+  // same row (stride 1) instead of striding across rows (stride size_n/8).
+  int col = blockIdx.y * 32 + threadIdx.x;
+  if (col >= size_n / 8) return;
+  (void)size_k;
+
+  int32_t val = qweight[blockIdx.x * (size_n / 8) + col];
+  int32_t zero = qzeros[blockIdx.x / group_size * (size_n / 8) + col];
  int32_t new_val = 0;

 #pragma unroll
@@ -58,7 +62,7 @@ __global__ void marlin_int4_fp8_preprocess_kernel_awq(
    zero >>= 4;
  }

-  output[(blockIdx.x * 32 + threadIdx.x) * size_n / 8 + blockIdx.y] = new_val;
+  output[blockIdx.x * (size_n / 8) + col] = new_val;
 }

 torch::stable::Tensor marlin_int4_fp8_preprocess(
@@ -102,7 +106,7 @@ torch::stable::Tensor marlin_int4_fp8_preprocess(
                    "qweight.size(0) % qzeros.size(0) != 0");
    STD_TORCH_CHECK(group_size % 8 == 0, "group_size % 8 != 0");

-    dim3 blocks(size_k / 32, size_n / 8);
+    dim3 blocks(size_k, (size_n / 8 + 31) / 32);
    marlin_int4_fp8_preprocess_kernel_awq<<<blocks, 32, 0, stream>>>(
        reinterpret_cast<const int32_t*>(qweight.const_data_ptr()),
        reinterpret_cast<int32_t*>(output.mutable_data_ptr()),
@@ -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 "
@@ -793,7 +793,7 @@ RUN --mount=type=cache,target=/opt/uv/cache \
 # Install FlashInfer JIT cache (requires CUDA-version-specific index URL)
 # https://docs.flashinfer.ai/installation.html
 # From versions.json: .flashinfer.version
-ARG FLASHINFER_VERSION=0.6.14
+ARG FLASHINFER_VERSION=0.6.15.post1
 RUN --mount=type=cache,target=/opt/uv/cache \
    uv pip install --system flashinfer-jit-cache==${FLASHINFER_VERSION} \
        --index-url https://flashinfer.ai/whl/cu$(echo $CUDA_VERSION | cut -d. -f1,2 | tr -d '.')
@@ -339,18 +339,17 @@ COPY --from=build_vllm ${COMMON_WORKDIR}/vllm/rust /rust
 COPY --from=build_vllm ${COMMON_WORKDIR}/vllm/rust-toolchain.toml /rust-toolchain.toml
 COPY --from=build_vllm ${COMMON_WORKDIR}/vllm/vllm/v1 /vllm_v1

-# RIXL/UCX build stages
-FROM base AS build_rixl
-ARG RIXL_BRANCH="39be1de8"
-ARG RIXL_REPO="https://github.com/ROCm/RIXL.git"
-ARG UCX_BRANCH="bfb51733"
+# NIXL/UCX build stages
+FROM base AS build_nixl
+ARG NIXL_BRANCH="231d56753047c989062a5cb2ac703a1ad761c7d2"
+ARG NIXL_REPO="https://github.com/ai-dynamo/nixl.git"
+ARG UCX_BRANCH="96e58a16039f6d7d213bc967b8069238742c5194"
 ARG UCX_REPO="https://github.com/openucx/ucx.git"
 ENV ROCM_PATH=/opt/rocm
 ENV UCX_HOME=/usr/local/ucx
-ENV RIXL_HOME=/usr/local/rixl
-ENV RIXL_BENCH_HOME=/usr/local/rixl_bench
+ENV NIXL_HOME=/usr/local/nixl

-# RIXL build system dependences and RDMA support
+# NIXL build system dependencies and RDMA support
 RUN apt-get -y update && apt-get -y install autoconf libtool pkg-config \
    libgrpc-dev \
    libgrpc++-dev \
@@ -368,7 +367,8 @@ RUN apt-get -y update && apt-get -y install autoconf libtool pkg-config \
    && rm -rf /var/lib/apt/lists/*

 RUN --mount=type=cache,target=/root/.cache/uv \
-    uv pip install --system meson auditwheel patchelf tomlkit
+    uv pip install --system meson meson-python pybind11 pyyaml types-PyYAML \
+        auditwheel build patchelf pytest tomlkit "setuptools>=80.9.0"

 RUN --mount=type=cache,target=/root/.cache/ccache \
    cd /usr/local/src && \
@@ -396,30 +396,50 @@ ENV PATH=/usr/local/ucx/bin:$PATH
 ENV LD_LIBRARY_PATH=${UCX_HOME}/lib:${LD_LIBRARY_PATH}

 RUN --mount=type=cache,target=/root/.cache/ccache \
-    git clone ${RIXL_REPO} /opt/rixl && \
-    cd /opt/rixl && \
-    git checkout ${RIXL_BRANCH} && \
+    git clone ${NIXL_REPO} /opt/nixl && \
+    cd /opt/nixl && \
+    git checkout ${NIXL_BRANCH} && \
    CC="ccache gcc" CXX="ccache g++" \
-    meson setup build --prefix=${RIXL_HOME} \
+    meson setup build --prefix=${NIXL_HOME} \
                     -Ducx_path=${UCX_HOME} \
-                     -Drocm_path=${ROCM_PATH} && \
+                     -Dwheel_variant=rocm \
+                     -Dbuild_tests=false \
+                     -Dbuild_examples=false && \
    cd build && \
    ninja -j$(nproc) && \
-    ninja install
+    ninja install && \
+    echo "${NIXL_HOME}/lib/$(uname -m)-linux-gnu" \
+        > /etc/ld.so.conf.d/nixl.conf && \
+    echo "${NIXL_HOME}/lib/$(uname -m)-linux-gnu/plugins" \
+        >> /etc/ld.so.conf.d/nixl.conf && \
+    ldconfig

-# Generate RIXL wheel
+# Generate the ROCm NIXL wheel. Upstream's generic wheel helper detects CUDA,
+# so configure the ROCm wheel variant directly through Meson.
 # Exclude libcore and libpull from auditwheel: transitive dependencies
 # that are not shipped in the wheel and vary across base images.
-RUN cd /opt/rixl && \
-    sed -i "s/--exclude 'libamdhip64\*'/--exclude 'libamdhip64*' --exclude 'libcore*' --exclude 'libpull*'/" \
-        contrib/build-wheel.sh && \
-    mkdir -p /app/install && \
-    _ucx_install_dir=${UCX_HOME} \
-    ./contrib/build-wheel.sh \
-        --output-dir /app/install \
-        --rocm-dir ${ROCM_PATH} \
+RUN cd /opt/nixl && \
+    ./contrib/tomlutil.py --wheel-name nixl-rocm pyproject.toml && \
+    CC="ccache gcc" CXX="ccache g++" \
+    uv build --wheel --no-build-isolation --out-dir /tmp/nixl_wheels \
+        --python ${PYTHON_VERSION} \
+        -Csetup-args=-Ducx_path=${UCX_HOME} \
+        -Csetup-args=-Dwheel_variant=rocm \
+        -Csetup-args=-Dbuild_tests=false \
+        -Csetup-args=-Dbuild_examples=false && \
+    mkdir -p /tmp/nixl_wheels/repaired /app/install && \
+    auditwheel repair \
+        --exclude 'libamdhip64*' \
+        --exclude 'libcore*' \
+        --exclude 'libpull*' \
+        /tmp/nixl_wheels/nixl_rocm*.whl \
+        --plat manylinux_2_34_$(uname -m) \
+        --wheel-dir /tmp/nixl_wheels/repaired && \
+    ./contrib/wheel_add_ucx_plugins.py \
        --ucx-plugins-dir ${UCX_HOME}/lib/ucx \
-        --nixl-plugins-dir ${RIXL_HOME}/lib/x86_64-linux-gnu/plugins
+        --nixl-plugins-dir ${NIXL_HOME}/lib/$(uname -m)-linux-gnu/plugins \
+        /tmp/nixl_wheels/repaired/*.whl && \
+    cp /tmp/nixl_wheels/repaired/*.whl /app/install

 # ROCShmem build stage - split from DeepEP so changing DEEPEP_BRANCH does not
 # invalidate the slow ROCShmem build.
@@ -660,10 +680,10 @@ RUN if [ "${DEEPEP_NIC}" = "cx7" ] || [ "${DEEPEP_NIC}" = "io" ]; then \
    ninja && ninja install && ldconfig && rm -rf /tmp/rdma-core; \
 fi

-# Install RIXL + DeepEP wheels.
-RUN --mount=type=bind,from=build_rixl,src=/app/install,target=/rixl_install \
+# Install NIXL + DeepEP wheels.
+RUN --mount=type=bind,from=build_nixl,src=/app/install,target=/nixl_install \
    --mount=type=bind,from=build_deepep,src=/app/deep_install,target=/deep_install \
-    uv pip install --system /rixl_install/*.whl /deep_install/*.whl
+    uv pip install --system /nixl_install/*.whl /deep_install/*.whl

 # Copy ROCShmem runtime libraries.
 COPY --from=build_rocshmem /opt/rocshmem /opt/rocshmem
@@ -724,6 +744,7 @@ ENV MIOPEN_DEBUG_CONV_GEMM=0
 # Use legacy IPC mode for HSA to avoid GPU memory pinning issues with UCX rocm_ipc.
 # See: https://github.com/ROCm/rocm-libraries/issues/6266
 ENV HSA_ENABLE_IPC_MODE_LEGACY=1
+ENV UCX_RMA_PPLN_ENABLE=y

 # ROCm profiler limits workaround.
 RUN echo "ROCTRACER_MAX_EVENTS=10000000" > ${COMMON_WORKDIR}/libkineto.conf
@@ -796,9 +817,9 @@ RUN --mount=type=bind,from=export_vllm,src=/,target=/install \
    && pip uninstall -y vllm \
    && uv pip install --system *.whl

-# Install RIXL wheel
-RUN --mount=type=bind,from=build_rixl,src=/app/install,target=/rixl_install \
-    uv pip install --system /rixl_install/*.whl
+# Install NIXL ROCm wheel
+RUN --mount=type=bind,from=build_nixl,src=/app/install,target=/nixl_install \
+    uv pip install --system /nixl_install/*.whl

 ARG COMMON_WORKDIR
 ARG BASE_IMAGE
@@ -813,6 +834,7 @@ COPY --from=export_vllm /docker ${COMMON_WORKDIR}/vllm/docker
 # Use legacy IPC mode for HSA to avoid GPU memory pinning issues with UCX rocm_ipc
 # See: https://github.com/ROCm/rocm-libraries/issues/6266
 ENV HSA_ENABLE_IPC_MODE_LEGACY=1
+ENV UCX_RMA_PPLN_ENABLE=y

 ENV TOKENIZERS_PARALLELISM=false

@@ -9,7 +9,7 @@ ARG PYTORCH_AUDIO_BRANCH="v2.9.0"
 ARG PYTORCH_AUDIO_REPO="https://github.com/pytorch/audio.git"
 ARG FA_BRANCH="0e60e394"
 ARG FA_REPO="https://github.com/Dao-AILab/flash-attention.git"
-ARG AITER_BRANCH="v0.1.16.post3"
+ARG AITER_BRANCH="v0.1.16.post5"
 ARG AITER_REPO="https://github.com/ROCm/aiter.git"
 ARG MORI_BRANCH="v1.1.0"
 ARG MORI_REPO="https://github.com/ROCm/mori.git"
@@ -86,6 +86,29 @@ RUN --mount=type=cache,target=/root/.cache/uv \
    mkdir -p /tmp/hf-xet/dist && \
    cp dist/*.whl /tmp/hf-xet/dist/

+# Build LLVM 20 from source for llvmlite (system repos ship LLVM 21 which
+# llvmlite v0.47 does not support; only SystemZ target is needed).
+FROM base AS llvm20-build
+ARG LLVM_VERSION=20.1.8
+WORKDIR /tmp
+RUN microdnf install -y ninja-build gcc gcc-c++ python3 xz && \
+    curl -LO https://github.com/llvm/llvm-project/releases/download/llvmorg-${LLVM_VERSION}/llvm-project-${LLVM_VERSION}.src.tar.xz && \
+    tar -xf llvm-project-${LLVM_VERSION}.src.tar.xz && \
+    cmake -G Ninja -S llvm-project-${LLVM_VERSION}.src/llvm -B build \
+        -DCMAKE_BUILD_TYPE=Release \
+        -DCMAKE_INSTALL_PREFIX=/opt/llvm20 \
+        -DLLVM_TARGETS_TO_BUILD="SystemZ" \
+        -DLLVM_ENABLE_RTTI=ON \
+        -DLLVM_BUILD_TOOLS=OFF \
+        -DLLVM_BUILD_UTILS=ON \
+        -DLLVM_BUILD_EXAMPLES=OFF \
+        -DLLVM_BUILD_TESTS=OFF \
+        -DLLVM_INCLUDE_TESTS=OFF \
+        -DLLVM_INCLUDE_EXAMPLES=OFF \
+        -DLLVM_INCLUDE_BENCHMARKS=OFF && \
+    ninja -C build install && \
+    rm -rf build llvm-project-${LLVM_VERSION}.src*
+
 # Build numba
 FROM python-install AS numba-builder

@@ -96,11 +119,13 @@ WORKDIR /tmp

 # Clone all required dependencies
 RUN --mount=type=cache,target=/root/.cache/uv \
+    --mount=type=bind,from=llvm20-build,source=/opt/llvm20,target=/opt/llvm20 \
    microdnf install ninja-build gcc gcc-c++ -y && \
    git clone --recursive https://github.com/numba/llvmlite.git -b v0.47.0 && \
    git clone --recursive https://github.com/numba/numba.git -b ${NUMBA_VERSION} && \
    cd llvmlite && \
    uv pip install 'cmake<4' 'setuptools<70' numpy && \
+    CMAKE_PREFIX_PATH=/opt/llvm20 LLVM_CONFIG=/opt/llvm20/bin/llvm-config \
    python setup.py bdist_wheel && \
    cd ../numba && \
    if ! grep '#include "dynamic_annotations.h"' numba/_dispatcher.cpp; then \
@@ -158,7 +183,6 @@ RUN --mount=type=cache,target=/root/.cache/uv \
     NUMBA_WHL_FILE=$(ls /tmp/numba-wheels/*.whl) && \
     OPENCV_WHL_FILE=$(ls /tmp/opencv-wheels/*.whl) && \
     uv pip install -v \
-        $ARROW_WHL_FILE \
        $VISION_WHL_FILE \
        $HF_XET_WHL_FILE \
        $LLVM_WHL_FILE \
@@ -59,7 +59,7 @@ variable "PYTORCH_ROCM_ARCH" {
 }

 # Pre-built CI base image (Tier 1). Per-PR builds pull this instead of
-# rebuilding RIXL/DeepEP/torchcodec from scratch. The ci_base stage in
+# rebuilding NIXL/DeepEP/torchcodec from scratch. The ci_base stage in
 # Dockerfile.rocm inherits from base, so CI_BASE_IMAGE only affects the test
 # stage and is irrelevant when building --target ci_base itself.
 variable "CI_BASE_IMAGE" {
@@ -75,7 +75,7 @@ variable "CI_MAX_JOBS" {
 # Upstream dependency commit pins -- extracted from Dockerfile.rocm by
 # ci-bake-rocm.sh at build time. Empty defaults are safe: the cache
 # functions produce no entries when the variable is empty.
-variable "RIXL_BRANCH" {
+variable "NIXL_BRANCH" {
  default = ""
 }

@@ -91,7 +91,7 @@ variable "DEEPEP_BRANCH" {
  default = ""
 }

-variable "RIXL_CACHE_KEY" {
+variable "NIXL_CACHE_KEY" {
  default = ""
 }

@@ -236,7 +236,7 @@ function "get_cache_to_rocm_rust" {
  ])
 }

-# Cache functions for upstream dependency stages (RIXL/UCX, ROCShmem, DeepEP).
+# Cache functions for upstream dependency stages (NIXL/UCX, ROCShmem, DeepEP).
 # These stages are pinned to specific upstream commit hashes, so cache keys use
 # those hashes rather than the Buildkite commit. This means the cache persists
 # across all vLLM commits as long as the upstream dependency pins don't change.
@@ -244,16 +244,16 @@ function "get_cache_to_rocm_rust" {
 function "get_cache_from_rocm_deps" {
  params = []
  result = compact([
-    RIXL_CACHE_KEY != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:rixl-rocm-${RIXL_CACHE_KEY}" : (RIXL_BRANCH != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:rixl-rocm-${RIXL_BRANCH}-ucx-${UCX_BRANCH}" : ""),
+    NIXL_CACHE_KEY != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:nixl-rocm-${NIXL_CACHE_KEY}" : (NIXL_BRANCH != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:nixl-rocm-${NIXL_BRANCH}-ucx-${UCX_BRANCH}" : ""),
    ROCSHMEM_CACHE_KEY != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:rocshmem-rocm-${ROCSHMEM_CACHE_KEY}" : (ROCSHMEM_BRANCH != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:rocshmem-rocm-${ROCSHMEM_BRANCH}" : ""),
    DEEPEP_CACHE_KEY != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:deepep-rocm-${DEEPEP_CACHE_KEY}" : (DEEPEP_BRANCH != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:deepep-rocm-${DEEPEP_BRANCH}-rocshmem-${ROCSHMEM_BRANCH}" : ""),
  ])
 }

-function "get_cache_to_rocm_rixl" {
+function "get_cache_to_rocm_nixl" {
  params = []
  result = compact([
-    RIXL_CACHE_KEY != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:rixl-rocm-${RIXL_CACHE_KEY},mode=min" : (RIXL_BRANCH != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:rixl-rocm-${RIXL_BRANCH}-ucx-${UCX_BRANCH},mode=min" : ""),
+    NIXL_CACHE_KEY != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:nixl-rocm-${NIXL_CACHE_KEY},mode=min" : (NIXL_BRANCH != "" ? "type=registry,ref=${DOCKERHUB_CACHE_REPO}:nixl-rocm-${NIXL_BRANCH}-ucx-${UCX_BRANCH},mode=min" : ""),
  ])
 }

@@ -372,11 +372,11 @@ variable "CI_BASE_IMAGE_TAG_STABLE" {
 # in the registry cache keyed by its upstream commit hash. When ci_base rebuilds
 # (e.g., requirements change), these stages are cache hits if their upstream
 # pins haven't changed -- saving ~35min of compilation.
-target "rixl-rocm-ci" {
+target "nixl-rocm-ci" {
  inherits   = ["_common-rocm", "_ci-rocm"]
-  target     = "build_rixl"
+  target     = "build_nixl"
  cache-from = get_cache_from_rocm_deps()
-  cache-to   = get_cache_to_rocm_rixl()
+  cache-to   = get_cache_to_rocm_nixl()
  output     = ["type=cacheonly"]
 }

@@ -396,7 +396,7 @@ target "deepep-rocm-ci" {
  output     = ["type=cacheonly"]
 }

-# Builds only the ci_base stage (RIXL, DeepEP, torchcodec, etc.)
+# Builds only the ci_base stage (NIXL, DeepEP, torchcodec, etc.)
 # Invoked by the ensure-ci-base step when the content hash of ci_base-affecting
 # files drifts from the remote image label. Per-PR builds then pull the result
 # as CI_BASE_IMAGE instead of rebuilding those slow layers on every commit.
@@ -412,7 +412,7 @@ target "ci-base-rocm-ci" {
      CI_BASE_IMAGE_TAG_CONTENT_EXTRA != "" ? "type=registry,ref=${CI_BASE_IMAGE_TAG_CONTENT_EXTRA}" : "",
      CI_BASE_IMAGE_TAG_STABLE != "" ? "type=registry,ref=${CI_BASE_IMAGE_TAG_STABLE}" : "",
    ]),
-    # Import upstream dependency caches so RIXL/ROCShmem/DeepEP stages
+    # Import upstream dependency caches so NIXL/ROCShmem/DeepEP stages
    # are cache hits even when ci_base itself needs rebuilding.
    get_cache_from_rocm_deps(),
  )
@@ -424,5 +424,5 @@ target "ci-base-rocm-ci" {
 # Group for ci_base builds -- exports dependency stage caches alongside the
 # ci_base image so future rebuilds can reuse them independently.
 group "ci-base-rocm-ci-with-deps" {
-  targets = ["rixl-rocm-ci", "rocshmem-rocm-ci", "deepep-rocm-ci", "ci-base-rocm-ci"]
+  targets = ["nixl-rocm-ci", "rocshmem-rocm-ci", "deepep-rocm-ci", "ci-base-rocm-ci"]
 }
@@ -53,7 +53,7 @@ variable "CI_BASE_IMAGE" {
 # Upstream dependency commit pins. Plain local bake builds use the Dockerfile
 # ARG defaults. ci-bake-rocm.sh resolves those defaults (plus any env
 # overrides) and writes a small HCL override before invoking CI targets.
-variable "RIXL_BRANCH" {
+variable "NIXL_BRANCH" {
  default = ""
 }

@@ -106,7 +106,7 @@ target "test-rocm" {
  output   = ["type=docker"]
 }

-# CI base image target - builds only the ci_base stage (RIXL, DeepEP,
+# CI base image target - builds only the ci_base stage (NIXL, DeepEP,
 # torchcodec, requirements, etc.). Used by the weekly scheduled build and
 # the auto-rebuild trigger when requirements change in a PR.
 target "ci-base-rocm" {
@@ -68,7 +68,7 @@
      "default": "true"
    },
    "FLASHINFER_VERSION": {
-      "default": "0.6.14"
+      "default": "0.6.15.post1"
    },
    "GDRCOPY_CUDA_VERSION": {
      "default": "12.8"
@@ -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`.
@@ -164,8 +164,8 @@ Priority is **1 = highest** (tried first).
 | `FLASHINFER` | XQA† | fp16, bf16 | `auto`, `float16`, `bfloat16`, `fp8`, `fp8_e4m3`, `fp8_e5m2` | 16, 32, 64, 128, 256, 512, 1024 | 64, 128, 256, 512 | ❌ | ❌ | ❌ | ✅ | Decoder | 9.0 |
 | `FLASHINFER` | trtllm-gen† | fp16, bf16 | `auto`, `float16`, `bfloat16`, `fp8`, `fp8_e4m3`, `fp8_e5m2`, `nvfp4` | 16, 32, 64, 128, 256, 512, 1024 | 64, 128, 256, 512 | ✅ | ✅ | ❌ | ✅ | Decoder | 10.x |
 | `FLASH_ATTN` | FA2* | fp16, bf16 | `auto`, `float16`, `bfloat16` | %16 | Any | ❌ | ✅ | ❌ | ✅ | All | ≥8.0 |
-| `FLASH_ATTN` | FA3* | fp16, bf16 | `auto`, `float16`, `bfloat16`, `fp8`, `fp8_e4m3`, `fp8_e5m2` | %16 | Any | ✅ | ✅ | ❌ | ✅ | All | 9.x |
-| `FLASH_ATTN` | FA4* | fp16, bf16 | `auto`, `float16`, `bfloat16` | %16 | Any | ✅ | ✅ | ❌ | ✅ | All | ≥10.0 |
+| `FLASH_ATTN` | FA3* | fp16, bf16 | `auto`, `float16`, `bfloat16`, `fp8`, `fp8_e4m3` | %16 | Any | ✅ | ✅ | ❌ | ✅ | All | 9.x |
+| `FLASH_ATTN` | FA4* | fp16, bf16 | `auto`, `float16`, `bfloat16`, `fp8`, `fp8_e4m3` | %16 | Any | ✅ | ✅ | ❌ | ✅ | All | ≥10.0 |
 | `FLASH_ATTN_DIFFKV` | | fp16, bf16 | `auto` | Any | Any | ❌ | ❌ | ❌ | ✅ | Decoder | Any |
 | `FLEX_ATTENTION` | | fp16, bf16, fp32 | `auto`, `float16`, `bfloat16` | %16 | Any | ❌ | ✅ | ✅ | ❌ | Decoder, Encoder Only | Any |
 | `HPC_ATTN` | | fp16, bf16 | `auto`, `bfloat16`, `fp8_e4m3` | 64 | 128 | ❌ | ❌ | ❌ | ❌ | Decoder | ≥9.0 |
@@ -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` | both | 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. With `TieringOffloadingSpec`, a CPU promotion is self-describing when a local request observes its primary-tier `HIT` before event translation; otherwise its stored event may retain the placeholder, while a later `HIT` can backfill metadata for removal. Pending-removal/re-promotion races and externally initiated promotions may also produce placeholders, and consumers must ignore removals for unknown hashes. 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.

@@ -13,11 +13,7 @@ Install the NIXL library: `uv pip install nixl`, as a quick start on Nvidia plat
 - Refer to [NIXL official repository](https://github.com/ai-dynamo/nixl) for more installation instructions
 - The specified required NIXL version can be found in [requirements/kv_connectors.txt](../../requirements/kv_connectors.txt) and other relevant config files

-For ROCm platform, the [ROCm docker file](../../docker/Dockerfile.rocm) includes RIXL and ucx already.
-
- Refer to [RIXL official repository](https://github.com/rocm/rixl) for more information
- The supportive libraries for RIXL can be found in [requirements/kv_connectors_rocm.txt](../../requirements/kv_connectors_rocm.txt)
- In the future we may remove RIXL from docker image file and users will be able to install from pre-compiled binary packages
+For ROCm, the [ROCm Dockerfile](../../docker/Dockerfile.rocm) builds NIXL and UCX with ROCm support from source.

 For non-cuda platform, please install nixl with ucx build from source, instructed as below.

@@ -315,7 +315,7 @@ vLLM CPU supports data parallel (DP), tensor parallel (TP) and pipeline parallel
 - vLLM CPU supports quantizations:
    - AWQ (x86 only)
    - GPTQ (x86 only)
-    - compressed-tensor INT8 W8A8 (x86, s390x)
+    - compressed-tensor INT8 W8A8 (x86 only)

 ### Why do I see `get_mempolicy: Operation not permitted` when running in Docker?

@@ -11,7 +11,7 @@ Currently, the CPU implementation for s390x architecture supports FP32, BF16 and
 - OS: `Linux`
 - SDK: `gcc/g++ >= 14.0.0` or later with Command Line Tools
 - Instruction Set Architecture (ISA): VXE support is required. Works with Z14 and above.
- Build install python packages: `torchvision`, `llvmlite`, `numba`, `pyarrow (for testing)`, `opencv-headless`
+- Build from source python packages (no pre-built s390x wheels): `torchvision`, `llvmlite`, `numba`, `opencv-python-headless`, `hf-xet`

 --8<-- [end:requirements]
 --8<-- [start:set-up-using-python]
@@ -28,13 +28,24 @@ Install the following packages from the package manager before building the vLLM

 ```bash
 dnf install -y \
-    which procps findutils tar vim git gcc-toolset-14 gcc-toolset-14-binutils gcc-toolset-14-libatomic-devel zlib-devel \
+    which procps findutils tar vim git patch xz ninja-build \
+    gcc-toolset-14 gcc-toolset-14-binutils gcc-toolset-14-libatomic-devel zlib-devel \
    libjpeg-turbo-devel libtiff-devel libpng-devel libwebp-devel freetype-devel harfbuzz-devel \
    openssl-devel openblas openblas-devel autoconf automake libtool cmake numpy libsndfile \
    clang llvm-devel llvm-static clang-devel
 ```

-Install rust>=1.80 which is needed for `outlines-core` and `uvloop` python packages installation.
+Build and install `numactl` from source:
+
+```bash
+curl -LO https://github.com/numactl/numactl/archive/refs/tags/v2.0.19.tar.gz
+tar -xvzf v2.0.19.tar.gz
+cd numactl-2.0.19
+./autogen.sh && ./configure && make && make install
+cd ..
+```
+
+Install rust>=1.80 which is needed for `outlines-core`, `uvloop`, and `hf-xet` python packages installation.

 ```bash
 curl https://sh.rustup.rs -sSf | sh -s -- -y && \
@@ -44,26 +55,79 @@ curl https://sh.rustup.rs -sSf | sh -s -- -y && \
 Execute the following commands to build and install vLLM from source.

 !!! tip
-    Please build the following dependencies, `torchvision`, `llvmlite`, `numba`, `llguidance`, `pyarrow`, `opencv-headless` from source before building vLLM.
+    Pre-built wheels are not available for s390x for the following packages. Build them from source before building vLLM: `torchvision`, `llvmlite`, `numba`, `opencv-python-headless`, `hf-xet`.
+    See `docker/Dockerfile.s390x` for exact versions and build commands used in each multi-stage build.
+
+!!! note "LLVM 20 required for llvmlite"
+    `llvmlite v0.47` requires LLVM 20, but UBI 9.6 repos ship LLVM 21 which is
+    not compatible. You must build LLVM 20 from source before building `llvmlite`:
+
+    ```bash
+    curl -LO https://github.com/llvm/llvm-project/releases/download/llvmorg-20.1.8/llvm-project-20.1.8.src.tar.xz
+    tar -xf llvm-project-20.1.8.src.tar.xz
+    cmake -G Ninja -S llvm-project-20.1.8.src/llvm -B llvm-build \
+        -DCMAKE_BUILD_TYPE=Release \
+        -DCMAKE_INSTALL_PREFIX=/opt/llvm20 \
+        -DLLVM_TARGETS_TO_BUILD="SystemZ" \
+        -DLLVM_ENABLE_RTTI=ON \
+        -DLLVM_BUILD_TOOLS=OFF \
+        -DLLVM_BUILD_UTILS=ON \
+        -DLLVM_BUILD_EXAMPLES=OFF \
+        -DLLVM_BUILD_TESTS=OFF \
+        -DLLVM_INCLUDE_TESTS=OFF \
+        -DLLVM_INCLUDE_EXAMPLES=OFF \
+        -DLLVM_INCLUDE_BENCHMARKS=OFF
+    ninja -C llvm-build install
+    ```
+
+    Then build `llvmlite` pointing to LLVM 20:
+
+    ```bash
+    CMAKE_PREFIX_PATH=/opt/llvm20 LLVM_CONFIG=/opt/llvm20/bin/llvm-config \
+        python setup.py bdist_wheel
+    ```

 ```bash
-    uv pip install -v \
-        -r requirements/build/cpu.txt \
-        -r requirements/cpu.txt \
-        --torch-backend cpu \
-        --index-strategy unsafe-best-match && \
-    VLLM_TARGET_DEVICE=cpu python setup.py bdist_wheel && \
-        uv pip install dist/*.whl
+uv pip install -v \
+    /path/to/torchvision.whl \
+    /path/to/llvmlite.whl \
+    /path/to/numba.whl \
+    /path/to/opencv_python_headless.whl \
+    /path/to/hf_xet.whl \
+    -r requirements/build/cpu.txt \
+    -r requirements/cpu.txt \
+    --torch-backend cpu \
+    --index-strategy unsafe-best-match && \
+VLLM_TARGET_DEVICE=cpu VLLM_CPU_MOE_PREPACK=0 python setup.py bdist_wheel && \
+    uv pip install dist/*.whl
 ```

 ??? console "pip"
    ```bash
-        pip install -v \
-            --extra-index-url https://download.pytorch.org/whl/cpu \
-            -r requirements/build/cpu.txt \
-            -r requirements/cpu.txt \
-        VLLM_TARGET_DEVICE=cpu python setup.py bdist_wheel && \
-            pip install dist/*.whl
+    pip install -v \
+        --extra-index-url https://download.pytorch.org/whl/cpu \
+        /path/to/torchvision.whl \
+        /path/to/llvmlite.whl \
+        /path/to/numba.whl \
+        /path/to/opencv_python_headless.whl \
+        /path/to/hf_xet.whl \
+        -r requirements/build/cpu.txt \
+        -r requirements/cpu.txt && \
+    VLLM_TARGET_DEVICE=cpu VLLM_CPU_MOE_PREPACK=0 python setup.py bdist_wheel && \
+        pip install dist/*.whl
+    ```
+
+!!! warning "Protobuf workaround for s390x"
+    The C++ protobuf extension crashes on s390x. After installation, set the
+    following environment variable and remove the C++ extensions:
+
+    ```bash
+    export PROTOCOL_BUFFERS_PYTHON_IMPLEMENTATION=python
+
+    # Remove C++ protobuf extensions that crash on s390x
+    SITE_PKGS=$(python -c "import site; print(site.getsitepackages()[0])")
+    rm -rf "$SITE_PKGS/google/_upb/"*.so \
+           "$SITE_PKGS/google/protobuf/pyext/"*.so 2>/dev/null || true
    ```

 --8<-- [end:build-wheel-from-source]
@@ -80,19 +144,20 @@ docker build -f docker/Dockerfile.s390x \

 # Launch OpenAI server
 docker run --rm \
-    --privileged true \
+    --security-opt seccomp=unconfined \
+    --cap-add SYS_NICE \
    --shm-size 4g \
    -p 8000:8000 \
    -e VLLM_CPU_KVCACHE_SPACE=<KV cache space> \
    -e VLLM_CPU_OMP_THREADS_BIND=<CPU cores for inference> \
    vllm-cpu-env \
    --model meta-llama/Llama-3.2-1B-Instruct \
-    --dtype float \
+    --dtype bfloat16 \
    other vLLM OpenAI server arguments
 ```

 !!! tip
-    An alternative of `--privileged true` is `--cap-add SYS_NICE --security-opt seccomp=unconfined`.
+    Alternatively, `--privileged=true` also works but is broader and not generally recommended.

 --8<-- [end:build-image-from-source]
 --8<-- [start:extra-information]
@@ -27,7 +27,7 @@ Currently, there are no pre-built XPU wheels.

 - First, install required [driver](https://dgpu-docs.intel.com/driver/installation.html#installing-gpu-drivers).
 - Second, install Python packages for vLLM XPU backend building (Intel OneAPI dependencies are installed automatically as part of `torch-xpu`, see [PyTorch XPU get started](https://docs.pytorch.org/docs/stable/notes/get_start_xpu.html)):
- Start from vllm-xpu-kernels v0.1.10, we recommend user upgrade driver to [compute runtime 26.18](https://github.com/intel/compute-runtime/releases/tag/26.14.37833.4) release, to avoid potential compatibility issue.
+- Start from vllm-xpu-kernels v0.1.10, we recommend user upgrade driver to [compute runtime 26.18](https://github.com/intel/compute-runtime/releases/tag/26.18.38308.1) release, to avoid potential compatibility issue.

 ```bash
 git clone https://github.com/vllm-project/vllm.git
@@ -58,7 +58,40 @@ VLLM_TARGET_DEVICE=xpu pip install --no-build-isolation -e . -v
 --8<-- [end:build-wheel-from-source]
 --8<-- [start:pre-built-images]

-Currently, we release prebuilt XPU images at docker [hub](https://hub.docker.com/r/intel/vllm/tags) based on vLLM released version. For more information, please refer release [note](https://github.com/intel/ai-containers/blob/main/vllm).
+vLLM offers official Docker images for deployment.
+The images can be used to run OpenAI compatible server and are available on Docker Hub as [vllm/vllm-openai-xpu](https://hub.docker.com/r/vllm/vllm-openai-xpu/tags).
+
+- `vllm/vllm-openai-xpu:latest` — stable release, available starting from v0.26.0
+- `vllm/vllm-openai-xpu:nightly` — preview build from the latest development branch, use this if you want the latest features and fixes
+
+```bash
+docker run --rm \
+    --network=host \
+    --device /dev/dri:/dev/dri \
+    -v /dev/dri/by-path:/dev/dri/by-path \
+    -v ~/.cache/huggingface:/root/.cache/huggingface \
+    --env "HF_TOKEN=$HF_TOKEN" \
+    --ipc=host \
+    --privileged \
+    vllm/vllm-openai-xpu:<tag> \
+    --model Qwen/Qwen3-0.6B
+```
+
+To use the docker image as base for development, you can launch it in interactive session through overriding the entrypoint.
+
+???+ console "Commands"
+    ```bash
+    docker run --rm -it \
+        --network=host \
+        --device /dev/dri:/dev/dri \
+        -v /dev/dri/by-path:/dev/dri/by-path \
+        -v ~/.cache/huggingface:/root/.cache/huggingface \
+        --env "HF_TOKEN=$HF_TOKEN" \
+        --ipc=host \
+        --privileged \
+        --entrypoint /bin/bash \
+        vllm/vllm-openai-xpu:<tag>
+    ```

 --8<-- [end:pre-built-images]
 --8<-- [start:build-image-from-source]
@@ -65,6 +65,15 @@ This guide will help you quickly get started with vLLM to perform:
    !!! tip
        A nightly Docker image is also available as [vllm/vllm-openai-rocm:nightly](https://hub.docker.com/r/vllm/vllm-openai-rocm/tags) for testing the latest development builds.

+=== "Intel GPU"
+
+    vLLM supports Intel GPUs through the XPU backend. Pre-built XPU wheels will be available soon.
+
+    Official Docker images for Intel GPUs are added to the vLLM release starting from v0.26.0. Nightly Docker image is also available as [vllm/vllm-openai-xpu:nightly](https://hub.docker.com/r/vllm/vllm-openai-xpu/tags).
+
+    !!! tip
+        For more detailed instructions, including building from source and Docker image setup, please refer to the [GPU installation guide](installation/gpu.md) and select the "Intel XPU" tab.
+
 === "Google TPU"

    To run vLLM on Google TPUs, you need to install the `vllm-tpu` package.
@@ -0,0 +1,3 @@
+// Reo.Dev documentation tracking
+// https://docs.reo.dev/integrations/input-sources/developer-insights/documentation
+!function(){var e,t,n;e="d5c4337961ef0ac",t=function(){Reo.init({clientID:"d5c4337961ef0ac", enableThirdPartyTracking: true})},(n=document.createElement("script")).src="https://static.reo.dev/"+e+"/reo.js",n.defer=!0,n.onload=t,document.head.appendChild(n)}();
@@ -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).
@@ -67,7 +67,7 @@ The Transcriptions API supports uploading audio files in various formats includi
 - `response_format`: Format of the response ("json", "text") (optional)
 - `temperature`: Sampling temperature between 0 and 1 (optional)

-For the complete list of supported parameters including sampling parameters and vLLM extensions, see the [protocol definitions](https://github.com/vllm-project/vllm/blob/main/vllm/entrypoints/openai/protocol.py#L2182).
+For the complete list of supported parameters including sampling parameters and vLLM extensions, see the [protocol definitions](https://github.com/vllm-project/vllm/blob/main/vllm/entrypoints/speech_to_text/transcription/protocol.py).

 **Response Format:**

@@ -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`.
@@ -160,3 +160,4 @@ extra_javascript:
  - https://unpkg.com/mathjax@3.2.2/es5/tex-mml-chtml.js
  - mkdocs/javascript/edit_and_feedback.js
  - mkdocs/javascript/slack_and_forum.js
+  - mkdocs/javascript/reo.js
@@ -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",
@@ -14,8 +14,8 @@ PyNvVideoCodec==2.0.4
 # flashinfer-cubin is not on PyPI since 0.6.14; setup.py excludes it from
 # install_requires so the published wheel does not carry an unresolvable pin
 --extra-index-url https://flashinfer.ai/whl/
-flashinfer-python==0.6.14
-flashinfer-cubin==0.6.14
+flashinfer-python==0.6.15.post1
+flashinfer-cubin==0.6.15.post1
 apache-tvm-ffi==0.1.10
 tilelang==0.1.9
 nvidia-cudnn-frontend>=1.19.1
@@ -26,7 +26,7 @@ fastsafetensors >= 0.3.2

 # QuACK and Cutlass DSL for FA4 (cute-DSL implementation)
 nvidia-cutlass-dsl[cu13]==4.6.0
-quack-kernels>=0.4.0  # Required for tml-fa4
+quack-kernels>=0.6.1  # Required for CUTLASS DSL 4.6 by MSA

 # Tokenspeed_MLA for faster mla with spec decode
 tokenspeed-mla==0.1.8; platform_system == "Linux"
@@ -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
@@ -12,4 +12,4 @@ ray[data]
 setuptools==78.1.0
 setuptools-rust>=1.9.0
 nixl==0.3.0
-tpu-inference==0.24.0
+tpu-inference==0.25.0
@@ -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"
@@ -3549,6 +3560,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"
@@ -4232,18 +4252,6 @@ dependencies = [
 "syn 2.0.117",
 ]

-[[package]]
-name = "subenum"
-version = "1.1.3"
-source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "ec3d08fe7078c57309d5c3d938e50eba95ba1d33b9c3a101a8465fc6861a5416"
-dependencies = [
- "heck",
- "proc-macro2",
- "quote",
- "syn 2.0.117",
-]
-
 [[package]]
 name = "subtle"
 version = "2.6.1"
@@ -4855,6 +4863,7 @@ dependencies = [
 "futures-core",
 "pin-project-lite",
 "tokio",
+ "tokio-util",
 ]

 [[package]]
@@ -4916,9 +4925,9 @@ dependencies = [

 [[package]]
 name = "tonic"
-version = "0.14.5"
+version = "0.14.6"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "fec7c61a0695dc1887c1b53952990f3ad2e3a31453e1f49f10e75424943a93ec"
+checksum = "ac2a5518c70fa84342385732db33fb3f44bc4cc748936eb5833d2df34d6445ef"
 dependencies = [
 "async-trait",
 "axum",
@@ -4945,9 +4954,9 @@ dependencies = [

 [[package]]
 name = "tonic-build"
-version = "0.14.5"
+version = "0.14.6"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "1882ac3bf5ef12877d7ed57aad87e75154c11931c2ba7e6cde5e22d63522c734"
+checksum = "c68f61875ac5293cf72e6c8cf0158086428c82c37229e98c840878f1706b0322"
 dependencies = [
 "prettyplease",
 "proc-macro2",
@@ -4956,10 +4965,23 @@ dependencies = [
 ]

 [[package]]
-name = "tonic-prost"
-version = "0.14.5"
+name = "tonic-health"
+version = "0.14.6"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "a55376a0bbaa4975a3f10d009ad763d8f4108f067c7c2e74f3001fb49778d309"
+checksum = "fcfab99db777fba2802f0dfa861d1628d1ae916fb199d29819941f139ae85082"
+dependencies = [
+ "prost",
+ "tokio",
+ "tokio-stream",
+ "tonic",
+ "tonic-prost",
+]
+
+[[package]]
+name = "tonic-prost"
+version = "0.14.6"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "50849f68853be452acf590cde0b146665b8d507b3b8af17261df47e02c209ea0"
 dependencies = [
 "bytes",
 "prost",
@@ -4968,9 +4990,9 @@ dependencies = [

 [[package]]
 name = "tonic-prost-build"
-version = "0.14.5"
+version = "0.14.6"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "f3144df636917574672e93d0f56d7edec49f90305749c668df5101751bb8f95a"
+checksum = "654e5643eff75d7f8c99197ce1440ed19a3474eada74c12bbac488b2cafdae27"
 dependencies = [
 "prettyplease",
 "proc-macro2",
@@ -5439,6 +5461,41 @@ 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",
+ "thiserror-ext",
+ "tiktoken-rs 0.9.1",
+ "tokenizers",
+ "tokio",
+ "tokio-stream",
+ "tracing",
+ "tracing-subscriber",
+ "url",
+ "uuid",
+]
+
 [[package]]
 name = "vllm-chat"
 version = "0.1.0"
@@ -5447,15 +5504,11 @@ dependencies = [
 "asynk-strim-attr",
 "bytes",
 "clap",
- "easy-ext",
 "expect-test",
 "futures",
 "half",
- "indexmap 2.13.0",
 "itertools 0.14.0",
 "llm-multimodal",
- "minijinja",
- "minijinja-contrib",
 "ndarray 0.17.2",
 "oss-harmony",
 "paste",
@@ -5466,17 +5519,16 @@ dependencies = [
 "serde_json",
 "serde_with",
 "serial_test",
- "strum",
- "subenum",
 "tempfile",
 "thiserror 2.0.18",
 "thiserror-ext",
- "time",
 "tokio",
 "tracing",
 "tracing-subscriber",
 "trait-set",
 "uuid",
+ "vllm-chat-renderer",
+ "vllm-chat-types",
 "vllm-engine-core-client",
 "vllm-llm",
 "vllm-parser",
@@ -5486,6 +5538,42 @@ dependencies = [
 "zeromq",
 ]

+[[package]]
+name = "vllm-chat-renderer"
+version = "0.1.0"
+dependencies = [
+ "anyhow",
+ "enum-as-inner",
+ "expect-test",
+ "indexmap 2.13.0",
+ "itertools 0.14.0",
+ "minijinja",
+ "minijinja-contrib",
+ "oss-harmony",
+ "serde",
+ "serde-json-fmt",
+ "serde_json",
+ "serde_with",
+ "strum",
+ "tempfile",
+ "thiserror 2.0.18",
+ "thiserror-ext",
+ "time",
+ "tracing",
+ "vllm-chat-types",
+ "vllm-tokenizer",
+]
+
+[[package]]
+name = "vllm-chat-types"
+version = "0.1.0"
+dependencies = [
+ "easy-ext",
+ "serde",
+ "serde_json",
+ "serde_with",
+]
+
 [[package]]
 name = "vllm-cmd"
 version = "0.1.0"
@@ -5507,6 +5595,7 @@ dependencies = [
 "tracing",
 "tracing-subscriber",
 "uuid",
+ "vllm-bench",
 "vllm-chat",
 "vllm-engine-core-client",
 "vllm-managed-engine",
@@ -5625,11 +5714,11 @@ dependencies = [
 "expect-test",
 "futures",
 "openai-protocol",
- "serde",
 "serde_json",
 "thiserror 2.0.18",
 "thiserror-ext",
 "tool-parser",
+ "vllm-chat-types",
 "vllm-tokenizer",
 "winnow",
 "xgrammar-structural-tag",
@@ -5676,6 +5765,7 @@ dependencies = [
 "tokio-stream",
 "tokio-util",
 "tonic",
+ "tonic-health",
 "tonic-prost",
 "tonic-prost-build",
 "tower",
@@ -5735,6 +5825,7 @@ dependencies = [
 "rustc-hash 1.1.0",
 "serde",
 "serde_json",
+ "serde_with",
 "tekken-rs",
 "tempfile",
 "thiserror 2.0.18",
@@ -6267,9 +6358,9 @@ checksum = "9edde0db4769d2dc68579893f2306b26c6ecfbe0ef499b013d731b7b9247e0b9"

 [[package]]
 name = "xgrammar-structural-tag"
-version = "0.1.0+xgrammar.0.2.2.4d145cc"
+version = "0.2.0+xgrammar.0.2.4.dd729e7"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "2436dea2393d55a3b188588aa300c5a8afe8f45a77da52c611fb4498a6c876e6"
+checksum = "d4d24c842efc3c24e9756aa426d530cbdac0980e49af223cb384e276e981ca0a"
 dependencies = [
 "auto_impl",
 "serde",
@@ -1,6 +1,9 @@
 [workspace]
 members = [
+    "src/bench",
    "src/chat",
+    "src/chat-renderer",
+    "src/chat-types",
    "src/cmd",
    "src/engine-core-client",
    "src/llm",
@@ -32,8 +35,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 +55,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 +78,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"
@@ -110,10 +120,11 @@ tokio = { version = "1.47.1", features = [
 tokio-openssl = "0.6"
 tokio-stream = "0.1"
 tokio-util = { version = "0.7.18", features = ["rt"] }
-tonic = "0.14.5"
-tonic-build = "0.14.5"
-tonic-prost = "0.14.5"
-tonic-prost-build = "0.14.5"
+tonic = "0.14.6"
+tonic-build = "0.14.6"
+tonic-health = "0.14.6"
+tonic-prost = "0.14.6"
+tonic-prost-build = "0.14.6"
 tool-parser = "1.2.0"
 tower = { version = "0.5.3", features = ["util"] }
 tower-http = { version = "0.6.8", features = ["cors", "trace"] }
@@ -121,9 +132,13 @@ 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-chat-renderer = { path = "src/chat-renderer" }
+vllm-chat-types = { path = "src/chat-types" }
 vllm-engine-core-client = { path = "src/engine-core-client" }
 vllm-llm = { path = "src/llm" }
 vllm-managed-engine = { path = "src/managed-engine" }
@@ -133,7 +148,7 @@ vllm-server = { path = "src/server" }
 vllm-text = { path = "src/text" }
 vllm-tokenizer = { path = "src/tokenizer" }
 winnow = { version = "1.0.2", features = ["simd"] }
-xgrammar-structural-tag = "0.1.0"
+xgrammar-structural-tag = "0.2.0"
 zeromq = { version = "0.6.0", default-features = false, features = [
    "tokio-runtime",
    "all-transport",
@@ -14,6 +14,10 @@ service Generate {
  rpc GenerateStream (GenerateRequest) returns (stream GenerateResponse) {}
 }

+service Control {
+  rpc Abort (AbortRequest) returns (AbortResponse) {}
+}
+
 // ======================================================================================
 // Generate Request
 // ======================================================================================
@@ -201,3 +205,12 @@ message TokenIds {
  repeated uint32 ids = 1;
 }

+// ======================================================================================
+// Control
+// ======================================================================================
+
+message AbortRequest {
+  repeated string request_ids = 1;
+}
+
+message AbortResponse {}
@@ -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,40 @@
+[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", "http2"] }
+rlimit.workspace = true
+rustc-hash.workspace = true
+serde = { workspace = true, features = ["rc"] }
+serde_json = { workspace = true, features = ["raw_value"] }
+thiserror.workspace = true
+thiserror-ext.workspace = true
+tiktoken-rs.workspace = true
+tokenizers.workspace = true
+tokio.workspace = true
+tokio-stream.workspace = true
+tracing.workspace = true
+tracing-subscriber.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,326 @@
+// 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 {
+                        tracing::warn!(
+                            backend = "vllm-rerank",
+                            inputs = list.len(),
+                            "rerank request has no documents"
+                        );
+                    }
+                    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() {
+                        tracing::warn!(
+                            backend = "vllm-rerank",
+                            dataset = "random",
+                            "rerank request has an empty query; use the random-rerank dataset"
+                        );
+                    }
+                    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());
+    }
+}
--- a/Show More
+++ b/Show More