Merge branch 'main' into wentao-skip-work-when-empty

refactor only inside nixl
Signed-off-by: yewentao256 <zhyanwentao@126.com>
2026-04-03 11:31:23 -04:00 · 2026-04-01 18:28:13 +00:00 · 2026-04-01 18:09:06 +00:00 · 2026-03-31 13:01:59 -04:00 · 2026-03-30 21:06:12 +00:00 · 2026-03-30 19:17:12 +00:00
44 changed files with 861 additions and 1843 deletions
@@ -5,6 +5,7 @@ steps:
    depends_on: []
    device: amd_cpu
    no_plugin: true
+    soft_fail: true
    commands:
    - >
      docker build
@@ -35,6 +35,23 @@ export PYTHONPATH=".."
 # Helper Functions
 ###############################################################################

+wait_for_clean_gpus() {
+  local timeout=${1:-300}
+  local start=$SECONDS
+  echo "--- Waiting for clean GPU state (timeout: ${timeout}s)"
+  while true; do
+    if grep -q clean /opt/amdgpu/etc/gpu_state; then
+      echo "GPUs state is \"clean\""
+      return
+    fi
+    if (( SECONDS - start >= timeout )); then
+      echo "Error: GPUs did not reach clean state within ${timeout}s" >&2
+      exit 1
+    fi
+    sleep 3
+  done
+}
+
 cleanup_docker() {
  # Get Docker's root directory
  docker_root=$(docker info -f '{{.DockerRootDir}}')
@@ -348,12 +365,19 @@ apply_rocm_test_overrides() {
 ###############################################################################

 # --- GPU initialization ---
+echo "--- Confirming Clean Initial State"
+wait_for_clean_gpus
+
 echo "--- ROCm info"
 rocminfo

 # --- Docker housekeeping ---
 cleanup_docker

+echo "--- Resetting GPUs"
+echo "reset" > /opt/amdgpu/etc/gpu_state
+wait_for_clean_gpus
+
 # --- Pull test image ---
 echo "--- Pulling container"
 image_name="rocm/vllm-ci:${BUILDKITE_COMMIT}"
@@ -224,6 +224,20 @@ steps:
  commands:
    - ./.buildkite/scripts/run-multi-node-test.sh /vllm-workspace/tests 2 2 $IMAGE_TAG "VLLM_TEST_SAME_HOST=0 torchrun --nnodes 2 --nproc-per-node=2 --rdzv_backend=c10d --rdzv_endpoint=192.168.10.10 distributed/test_same_node.py | grep 'Same node test passed' && NUM_NODES=2 torchrun --nnodes 2 --nproc-per-node=2 --rdzv_backend=c10d --rdzv_endpoint=192.168.10.10 distributed/test_node_count.py | grep 'Node count test passed' && python3 ../examples/offline_inference/data_parallel.py -dp=2 -tp=1 --dp-num-nodes=2 --dp-node-rank=0 --dp-master-addr=192.168.10.10 --dp-master-port=12345 --enforce-eager --trust-remote-code && VLLM_MULTI_NODE=1 pytest -v -s distributed/test_multi_node_assignment.py && VLLM_MULTI_NODE=1 pytest -v -s distributed/test_pipeline_parallel.py" "VLLM_TEST_SAME_HOST=0 torchrun --nnodes 2 --nproc-per-node=2 --rdzv_backend=c10d --rdzv_endpoint=192.168.10.10 distributed/test_same_node.py | grep 'Same node test passed' && NUM_NODES=2 torchrun --nnodes 2 --nproc-per-node=2 --rdzv_backend=c10d --rdzv_endpoint=192.168.10.10 distributed/test_node_count.py | grep 'Node count test passed' && python3 ../examples/offline_inference/data_parallel.py -dp=2 -tp=1 --dp-num-nodes=2 --dp-node-rank=1 --dp-master-addr=192.168.10.10 --dp-master-port=12345 --enforce-eager --trust-remote-code"

+- label: MessageQueue TCP Multi-Node (2 GPUs)
+  timeout_in_minutes: 10
+  working_dir: "/vllm-workspace/tests"
+  num_devices: 1
+  num_nodes: 2
+  no_plugin: true
+  optional: true
+  source_file_dependencies:
+  - vllm/distributed/device_communicators/shm_broadcast.py
+  - vllm/distributed/parallel_state.py
+  - tests/distributed/test_mq_tcp_multinode.py
+  commands:
+    - ./.buildkite/scripts/run-multi-node-test.sh /vllm-workspace/tests 2 1 $IMAGE_TAG "torchrun --nnodes 2 --nproc-per-node=1 --rdzv_backend=c10d --rdzv_endpoint=192.168.10.10 distributed/test_mq_tcp_multinode.py" "torchrun --nnodes 2 --nproc-per-node=1 --rdzv_backend=c10d --rdzv_endpoint=192.168.10.10 distributed/test_mq_tcp_multinode.py"
+
 - label: Distributed NixlConnector PD accuracy (4 GPUs)
  timeout_in_minutes: 30
  working_dir: "/vllm-workspace/tests"
@@ -78,6 +78,7 @@ steps:
    - TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_async_llm_dp.py -k "not ray"
    - TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_eagle_dp.py

+# These require fix https://github.com/vllm-project/vllm/pull/36280
 - label: Model Runner V2 Pipeline Parallelism (4 GPUs)
  timeout_in_minutes: 60
  working_dir: "/vllm-workspace/tests"
@@ -91,9 +91,9 @@ void swap_blocks_batch(const torch::Tensor& src_ptrs,

  if (n == 0) return;

-  int64_t* src_data = src_ptrs.mutable_data_ptr<int64_t>();
-  int64_t* dst_data = dst_ptrs.mutable_data_ptr<int64_t>();
-  int64_t* size_data = sizes.mutable_data_ptr<int64_t>();
+  const int64_t* src_data = src_ptrs.data_ptr<int64_t>();
+  const int64_t* dst_data = dst_ptrs.data_ptr<int64_t>();
+  const int64_t* size_data = sizes.data_ptr<int64_t>();

  const cudaStream_t stream = at::cuda::getCurrentCUDAStream();

@@ -107,24 +107,15 @@ void swap_blocks_batch(const torch::Tensor& src_ptrs,
  CUmemcpyAttributes attr = {};
  attr.srcAccessOrder = CU_MEMCPY_SRC_ACCESS_ORDER_STREAM;
  size_t attrs_idx = 0;
-  #if defined(CUDA_VERSION) && CUDA_VERSION >= 13000
-  CUresult result = cuMemcpyBatchAsync(
-      reinterpret_cast<CUdeviceptr*>(dst_data),
-      reinterpret_cast<CUdeviceptr*>(src_data),
-      reinterpret_cast<size_t*>(size_data), static_cast<size_t>(n), &attr,
-      &attrs_idx, 1, static_cast<CUstream>(stream));
-  TORCH_CHECK(result == CUDA_SUCCESS, "cuMemcpyBatchAsync failed with error ",
-              result);
-  #else
  size_t fail_idx = 0;
  CUresult result = cuMemcpyBatchAsync(
-      reinterpret_cast<CUdeviceptr*>(dst_data),
-      reinterpret_cast<CUdeviceptr*>(src_data),
-      reinterpret_cast<size_t*>(size_data), static_cast<size_t>(n), &attr,
-      &attrs_idx, 1, &fail_idx, static_cast<CUstream>(stream));
+      reinterpret_cast<CUdeviceptr*>(const_cast<int64_t*>(dst_data)),
+      reinterpret_cast<CUdeviceptr*>(const_cast<int64_t*>(src_data)),
+      reinterpret_cast<size_t*>(const_cast<int64_t*>(size_data)),
+      static_cast<size_t>(n), &attr, &attrs_idx, 1, &fail_idx,
+      static_cast<CUstream>(stream));
  TORCH_CHECK(result == CUDA_SUCCESS, "cuMemcpyBatchAsync failed at index ",
              fail_idx, " with error ", result);
-  #endif
 #else
  // Fallback for CUDA < 12.8 and ROCm: individual async copies.
  // cudaMemcpyDefault lets the driver infer direction from pointer types.
@@ -685,9 +685,6 @@ RUN --mount=type=bind,from=build,src=/workspace/dist,target=/vllm-workspace/dist
        --extra-index-url ${PYTORCH_CUDA_INDEX_BASE_URL}/cu$(echo $CUDA_VERSION | cut -d. -f1,2 | tr -d '.'); \
    fi

-RUN --mount=type=cache,target=/root/.cache/uv \
-    uv pip install --system "transformers==5.4.0"
-
 RUN --mount=type=cache,target=/root/.cache/uv \
 . /etc/environment && \
 uv pip list
@@ -390,21 +390,20 @@ ENV MIOPEN_DEBUG_CONV_GEMM=0
 RUN mkdir src && mv vllm src/vllm

 # This is a workaround to ensure pytest exits with the correct status code in CI tests.
-RUN printf '%s\n' \
-    'import os' \
-    '' \
-    '_exit_code = 1' \
-    '' \
-    'def pytest_sessionfinish(session, exitstatus):' \
-    '    global _exit_code' \
-    '    _exit_code = int(exitstatus)' \
-    '' \
-    'def pytest_unconfigure(config):' \
-    '    import sys' \
-    '    sys.stdout.flush()' \
-    '    sys.stderr.flush()' \
-    '    os._exit(_exit_code)' \
-    > /vllm-workspace/conftest.py
+RUN cat << 'EOF' > /vllm-workspace/conftest.py
+import os
+
+_exit_code = 1
+
+def pytest_sessionfinish(session, exitstatus):
+    global _exit_code
+    _exit_code = int(exitstatus)
+
+def pytest_unconfigure(config):
+    sys.stdout.flush()
+    sys.stderr.flush()
+    os._exit(_exit_code)
+EOF

 # -----------------------
 # Final vLLM image
@@ -1,5 +1,3 @@
-r common.txt
-
 # testing
 pytest
 tensorizer==2.10.1
@@ -15,7 +15,6 @@ aiohappyeyeballs==2.6.1
 aiohttp==3.13.3
    # via
    #   -c requirements/common.txt
-    #   -r requirements/common.txt
    #   aiohttp-cors
    #   fsspec
    #   gpt-oss
@@ -39,31 +38,20 @@ annotated-doc==0.0.4
    #   typer
 annotated-types==0.7.0
    # via pydantic
-anthropic==0.89.0
-    # via
-    #   -c requirements/common.txt
-    #   -r requirements/common.txt
 antlr4-python3-runtime==4.9.3
    # via
    #   hydra-core
    #   omegaconf
-anyio==4.13.0
+anyio==4.6.2.post1
    # via
-    #   anthropic
    #   httpx
-    #   mcp
-    #   openai
-    #   sse-starlette
    #   starlette
-    #   watchfiles
 arctic-inference==0.1.1
    # via -r requirements/rocm-test.in
 argcomplete==3.6.3
    # via datamodel-code-generator
 arrow==1.4.0
    # via isoduration
-astor==0.8.1
-    # via depyf
 attrs==26.1.0
    # via
    #   aiohttp
@@ -95,8 +83,6 @@ bitsandbytes==0.49.2
    #   lightning
 black==26.3.1
    # via datamodel-code-generator
-blake3==1.0.8
-    # via -r requirements/common.txt
 blobfile==3.0.0
    # via -r requirements/rocm-test.in
 bm25s==0.2.13
@@ -113,10 +99,6 @@ bounded-pool-executor==0.0.3
    # via pqdm
 buildkite-test-collector==0.1.9
    # via -r requirements/rocm-test.in
-cachetools==7.0.5
-    # via -r requirements/common.txt
-cbor2==5.9.0
-    # via -r requirements/common.txt
 certifi==2026.2.25
    # via
    #   fiona
@@ -150,7 +132,6 @@ click==8.3.1
    #   nltk
    #   rasterio
    #   ray
-    #   rich-toolkit
    #   schemathesis
    #   typer
    #   uvicorn
@@ -161,8 +142,6 @@ cligj==0.7.2
    # via
    #   fiona
    #   rasterio
-cloudpickle==3.1.2
-    # via -r requirements/common.txt
 colorama==0.4.6
    # via
    #   perceptron
@@ -172,10 +151,6 @@ colorful==0.5.8
    # via ray
 colorlog==6.10.1
    # via optuna
-compressed-tensors==0.14.0.1
-    # via
-    #   -c requirements/common.txt
-    #   -r requirements/common.txt
 contourpy==1.3.3
    # via matplotlib
 coverage==7.13.5
@@ -207,42 +182,24 @@ decorator==5.2.1
    # via librosa
 decord==0.6.0
    # via -r requirements/rocm-test.in
-depyf==0.20.0
-    # via
-    #   -c requirements/common.txt
-    #   -r requirements/common.txt
 diffusers==0.37.0
    # via terratorch
 dill==0.3.8
    # via
    #   datasets
-    #   depyf
    #   evaluate
    #   lm-eval
    #   multiprocess
-diskcache==5.6.3
-    # via
-    #   -c requirements/common.txt
-    #   -r requirements/common.txt
 distlib==0.4.0
    # via virtualenv
-distro==1.9.0
-    # via
-    #   anthropic
-    #   openai
-dnspython==2.8.0
-    # via email-validator
 docker==7.1.0
    # via gpt-oss
 docopt==0.6.2
    # via num2words
 docstring-parser==0.17.0
-    # via
-    #   anthropic
-    #   jsonargparse
+    # via jsonargparse
 einops==0.8.2
    # via
-    #   -r requirements/common.txt
    #   -r requirements/rocm-test.in
    #   encodec
    #   terratorch
@@ -251,10 +208,6 @@ einops==0.8.2
    #   vocos
 einx==0.4.2
    # via vector-quantize-pytorch
-email-validator==2.3.0
-    # via
-    #   fastapi
-    #   pydantic
 encodec==0.1.1
    # via vocos
 et-xmlfile==2.0.0
@@ -264,15 +217,7 @@ evaluate==0.4.6
 fastapi==0.135.2
    # via
    #   -c requirements/common.txt
-    #   -r requirements/common.txt
    #   gpt-oss
-    #   model-hosting-container-standards
-fastapi-cli==0.0.24
-    # via fastapi
-fastapi-cloud-cli==0.15.1
-    # via fastapi-cli
-fastar==0.9.0
-    # via fastapi-cloud-cli
 fastparquet==2026.3.0
    # via genai-perf
 fastsafetensors==0.2.2
@@ -280,7 +225,6 @@ fastsafetensors==0.2.2
 filelock==3.25.2
    # via
    #   -c requirements/common.txt
-    #   -r requirements/common.txt
    #   blobfile
    #   datasets
    #   diffusers
@@ -320,10 +264,6 @@ genson==1.3.0
    # via datamodel-code-generator
 geopandas==1.1.3
    # via terratorch
-gguf==0.18.0
-    # via
-    #   -c requirements/common.txt
-    #   -r requirements/common.txt
 gitdb==4.0.12
    # via gitpython
 gitpython==3.1.46
@@ -350,10 +290,7 @@ google-crc32c==1.8.0
 google-resumable-media==2.8.0
    # via google-cloud-storage
 googleapis-common-protos==1.73.0
-    # via
-    #   google-api-core
-    #   opentelemetry-exporter-otlp-proto-grpc
-    #   opentelemetry-exporter-otlp-proto-http
+    # via google-api-core
 gpt-oss==0.0.8
    # via -r requirements/rocm-test.in
 graphql-core==3.2.8
@@ -365,7 +302,6 @@ grpcio==1.78.0
    #   -c requirements/rocm.txt
    #   -r requirements/rocm-test.in
    #   grpcio-reflection
-    #   opentelemetry-exporter-otlp-proto-grpc
    #   ray
    #   tensorboard
 grpcio-reflection==1.78.0
@@ -392,22 +328,12 @@ html2text==2025.4.15
    # via gpt-oss
 httpcore==1.0.9
    # via httpx
-httptools==0.7.1
-    # via uvicorn
 httpx==0.27.2
    # via
    #   -r requirements/rocm-test.in
-    #   anthropic
    #   diffusers
-    #   fastapi
-    #   fastapi-cloud-cli
-    #   mcp
-    #   model-hosting-container-standards
-    #   openai
    #   perceptron
    #   schemathesis
-httpx-sse==0.4.3
-    # via mcp
 huggingface-hub==0.36.2
    # via
    #   -r requirements/rocm-test.in
@@ -444,13 +370,10 @@ hypothesis-jsonschema==0.23.1
 idna==3.11
    # via
    #   anyio
-    #   email-validator
    #   httpx
    #   jsonschema
    #   requests
    #   yarl
-ijson==3.5.0
-    # via -r requirements/common.txt
 imagehash==4.3.2
    # via -r requirements/rocm-test.in
 imageio==2.37.3
@@ -467,8 +390,6 @@ iniconfig==2.3.0
    # via pytest
 instanttensor==0.1.6
    # via -r requirements/rocm-test.in
-interegular==0.3.3
-    # via lm-format-enforcer
 isodate==0.7.2
    # via azure-storage-blob
 isoduration==20.11.0
@@ -478,21 +399,15 @@ isort==8.0.1
 jinja2==3.1.6
    # via
    #   datamodel-code-generator
-    #   fastapi
    #   genai-perf
    #   lm-eval
    #   torch
-jiter==0.13.0
-    # via
-    #   anthropic
-    #   openai
 jiwer==4.0.0
    # via -r requirements/rocm-test.in
 jmespath==1.1.0
    # via
    #   boto3
    #   botocore
-    #   model-hosting-container-standards
 joblib==1.5.3
    # via
    #   librosa
@@ -511,7 +426,6 @@ jsonpointer==3.1.0
 jsonschema==4.26.0
    # via
    #   hypothesis-jsonschema
-    #   mcp
    #   mistral-common
    #   ray
    #   schemathesis
@@ -529,10 +443,6 @@ kornia==0.8.2
    # via torchgeo
 kornia-rs==0.1.10
    # via kornia
-lark==1.2.2
-    # via
-    #   -c requirements/common.txt
-    #   -r requirements/common.txt
 lazy-loader==0.4
    # via
    #   librosa
@@ -556,24 +466,14 @@ lightning-utilities==0.15.3
    #   lightning
    #   pytorch-lightning
    #   torchmetrics
-llguidance==1.3.0
-    # via
-    #   -c requirements/common.txt
-    #   -r requirements/common.txt
 llvmlite==0.44.0
    # via numba
 lm-eval==0.4.11
    # via -r requirements/rocm-test.in
-lm-format-enforcer==0.11.3
-    # via
-    #   -c requirements/common.txt
-    #   -r requirements/common.txt
 logistro==2.0.1
    # via
    #   choreographer
    #   kaleido
-loguru==0.7.3
-    # via compressed-tensors
 lxml==6.0.2
    # via
    #   blobfile
@@ -600,19 +500,12 @@ mbstrdecoder==1.1.4
    #   dataproperty
    #   pytablewriter
    #   typepy
-mcp==1.27.0
-    # via -r requirements/common.txt
 mdurl==0.1.2
    # via markdown-it-py
 mistral-common==1.10.0
    # via
    #   -c requirements/common.txt
-    #   -r requirements/common.txt
    #   -r requirements/rocm-test.in
-model-hosting-container-standards==0.1.14
-    # via
-    #   -c requirements/common.txt
-    #   -r requirements/common.txt
 more-itertools==10.8.0
    # via
    #   inflect
@@ -629,8 +522,6 @@ msgpack==1.1.2
    # via
    #   librosa
    #   ray
-msgspec==0.20.0
-    # via -r requirements/common.txt
 mteb==2.11.5
    # via -r requirements/rocm-test.in
 multidict==6.7.1
@@ -650,8 +541,6 @@ networkx==3.6.1
    # via
    #   scikit-image
    #   torch
-ninja==1.13.0
-    # via -r requirements/common.txt
 nltk==3.9.3
    # via rouge-score
 num2words==0.5.14
@@ -666,7 +555,6 @@ numkong==7.1.1
    # via albucore
 numpy==2.2.6
    # via
-    #   -r requirements/common.txt
    #   -r requirements/rocm-test.in
    #   accelerate
    #   albucore
@@ -684,7 +572,6 @@ numpy==2.2.6
    #   fastparquet
    #   genai-perf
    #   geopandas
-    #   gguf
    #   h5py
    #   imagehash
    #   imageio
@@ -733,21 +620,15 @@ numpy==2.2.6
    #   tritonclient
    #   vocos
    #   xarray
-    #   xgrammar
 omegaconf==2.3.0
    # via
    #   hydra-core
    #   lightning
 open-clip-torch==2.32.0
    # via -r requirements/rocm-test.in
-openai==2.30.0
-    # via
-    #   -c requirements/common.txt
-    #   -r requirements/common.txt
 openai-harmony==0.0.8
    # via
    #   -c requirements/common.txt
-    #   -r requirements/common.txt
    #   gpt-oss
 opencensus==0.11.4
    # via ray
@@ -756,7 +637,6 @@ opencensus-context==0.1.3
 opencv-python-headless==4.13.0.92
    # via
    #   -c requirements/common.txt
-    #   -r requirements/common.txt
    #   -r requirements/rocm-test.in
    #   albumentations
    #   mistral-common
@@ -765,59 +645,26 @@ openpyxl==3.1.5
 opentelemetry-api==1.40.0
    # via
    #   -c requirements/common.txt
-    #   -r requirements/common.txt
-    #   opentelemetry-exporter-otlp-proto-grpc
-    #   opentelemetry-exporter-otlp-proto-http
    #   opentelemetry-exporter-prometheus
    #   opentelemetry-sdk
    #   opentelemetry-semantic-conventions
-opentelemetry-exporter-otlp==1.40.0
-    # via
-    #   -c requirements/common.txt
-    #   -r requirements/common.txt
-opentelemetry-exporter-otlp-proto-common==1.40.0
-    # via
-    #   opentelemetry-exporter-otlp-proto-grpc
-    #   opentelemetry-exporter-otlp-proto-http
-opentelemetry-exporter-otlp-proto-grpc==1.40.0
-    # via opentelemetry-exporter-otlp
-opentelemetry-exporter-otlp-proto-http==1.40.0
-    # via opentelemetry-exporter-otlp
 opentelemetry-exporter-prometheus==0.61b0
    # via ray
 opentelemetry-proto==1.40.0
-    # via
-    #   opentelemetry-exporter-otlp-proto-common
-    #   opentelemetry-exporter-otlp-proto-grpc
-    #   opentelemetry-exporter-otlp-proto-http
-    #   ray
+    # via ray
 opentelemetry-sdk==1.40.0
    # via
    #   -c requirements/common.txt
-    #   -r requirements/common.txt
-    #   opentelemetry-exporter-otlp-proto-grpc
-    #   opentelemetry-exporter-otlp-proto-http
    #   opentelemetry-exporter-prometheus
-    #   opentelemetry-semantic-conventions-ai
    #   ray
 opentelemetry-semantic-conventions==0.61b0
-    # via
-    #   opentelemetry-sdk
-    #   opentelemetry-semantic-conventions-ai
-opentelemetry-semantic-conventions-ai==0.5.1
-    # via
-    #   -c requirements/common.txt
-    #   -r requirements/common.txt
+    # via opentelemetry-sdk
 optuna==3.6.1
    # via genai-perf
 orjson==3.11.7
    # via
    #   genai-perf
    #   kaleido
-outlines-core==0.2.11
-    # via
-    #   -c requirements/common.txt
-    #   -r requirements/common.txt
 packaging==26.0
    # via
    #   -c requirements/rocm.txt
@@ -835,7 +682,6 @@ packaging==26.0
    #   lazy-loader
    #   lightning
    #   lightning-utilities
-    #   lm-format-enforcer
    #   matplotlib
    #   optuna
    #   peft
@@ -867,8 +713,6 @@ pandas==3.0.1
    #   tacoreader
    #   torchgeo
    #   xarray
-partial-json-parser==0.2.1.1.post7
-    # via -r requirements/common.txt
 pathspec==1.0.4
    # via black
 pathvalidate==3.3.1
@@ -883,7 +727,6 @@ perf-analyzer==0.1.0
    # via genai-perf
 pillow==12.1.1
    # via
-    #   -r requirements/common.txt
    #   diffusers
    #   genai-perf
    #   imagehash
@@ -925,14 +768,8 @@ pqdm==0.2.0
 prometheus-client==0.24.1
    # via
    #   -c requirements/common.txt
-    #   -r requirements/common.txt
    #   opentelemetry-exporter-prometheus
-    #   prometheus-fastapi-instrumentator
    #   ray
-prometheus-fastapi-instrumentator==7.1.0
-    # via
-    #   -c requirements/common.txt
-    #   -r requirements/common.txt
 propcache==0.4.1
    # via
    #   aiohttp
@@ -942,7 +779,6 @@ proto-plus==1.27.1
 protobuf==6.33.6
    # via
    #   -c requirements/common.txt
-    #   -r requirements/common.txt
    #   google-api-core
    #   googleapis-common-protos
    #   grpcio-reflection
@@ -955,14 +791,11 @@ protobuf==6.33.6
    #   wandb
 psutil==7.2.2
    # via
-    #   -r requirements/common.txt
    #   accelerate
    #   peft
    #   tensorizer
 py==1.11.0
    # via pytest-forked
-py-cpuinfo==9.0.0
-    # via -r requirements/common.txt
 py-spy==0.4.1
    # via ray
 pyarrow==23.0.1
@@ -975,8 +808,6 @@ pyasn1==0.6.3
    # via pyasn1-modules
 pyasn1-modules==0.4.2
    # via google-auth
-pybase64==1.4.3
-    # via -r requirements/common.txt
 pycocotools==2.0.11
    # via terratorch
 pycountry==26.2.16
@@ -988,44 +819,26 @@ pycryptodomex==3.23.0
 pydantic==2.12.5
    # via
    #   -c requirements/common.txt
-    #   -r requirements/common.txt
    #   -r requirements/rocm-test.in
    #   albumentations
-    #   anthropic
-    #   compressed-tensors
    #   datamodel-code-generator
    #   fastapi
-    #   fastapi-cloud-cli
    #   gpt-oss
    #   lightly
-    #   lm-format-enforcer
-    #   mcp
    #   mistral-common
-    #   model-hosting-container-standards
    #   mteb
-    #   openai
    #   openai-harmony
    #   pydantic-extra-types
-    #   pydantic-settings
    #   ray
    #   wandb
-    #   xgrammar
 pydantic-core==2.41.5
    # via pydantic
 pydantic-extra-types==2.11.1
-    # via
-    #   fastapi
-    #   mistral-common
-pydantic-settings==2.13.1
-    # via
-    #   fastapi
-    #   mcp
+    # via mistral-common
 pygments==2.19.2
    # via rich
 pyjwt==2.12.1
-    # via
-    #   mcp
-    #   msal
+    # via msal
 pyogrio==0.12.1
    # via geopandas
 pyparsing==3.3.2
@@ -1085,16 +898,6 @@ python-dateutil==2.9.0.post0
    #   typepy
 python-discovery==1.2.0
    # via virtualenv
-python-dotenv==1.2.2
-    # via
-    #   pydantic-settings
-    #   uvicorn
-python-json-logger==4.1.0
-    # via -r requirements/common.txt
-python-multipart==0.0.22
-    # via
-    #   fastapi
-    #   mcp
 python-rapidjson==1.23
    # via tritonclient
 pytokens==0.4.1
@@ -1111,17 +914,14 @@ pywavelets==1.9.0
    # via imagehash
 pyyaml==6.0.3
    # via
-    #   -r requirements/common.txt
    #   accelerate
    #   albumentations
    #   datamodel-code-generator
    #   datasets
    #   genai-perf
-    #   gguf
    #   huggingface-hub
    #   jsonargparse
    #   lightning
-    #   lm-format-enforcer
    #   omegaconf
    #   optuna
    #   peft
@@ -1131,13 +931,8 @@ pyyaml==6.0.3
    #   schemathesis
    #   timm
    #   transformers
-    #   uvicorn
    #   vocos
    #   wandb
-pyzmq==27.1.0
-    # via
-    #   -c requirements/common.txt
-    #   -r requirements/common.txt
 rapidfuzz==3.12.1
    # via
    #   -r requirements/rocm-test.in
@@ -1157,7 +952,6 @@ referencing==0.37.0
    #   jsonschema-specifications
 regex==2026.2.28
    # via
-    #   -r requirements/common.txt
    #   diffusers
    #   nltk
    #   open-clip-torch
@@ -1167,14 +961,12 @@ regex==2026.2.28
 requests==2.32.5
    # via
    #   -c requirements/common.txt
-    #   -r requirements/common.txt
    #   azure-core
    #   buildkite-test-collector
    #   datasets
    #   diffusers
    #   docker
    #   evaluate
-    #   gguf
    #   google-api-core
    #   google-cloud-storage
    #   gpt-oss
@@ -1184,7 +976,6 @@ requests==2.32.5
    #   mistral-common
    #   msal
    #   mteb
-    #   opentelemetry-exporter-otlp-proto-http
    #   pooch
    #   ray
    #   responses
@@ -1208,15 +999,8 @@ rich==14.3.3
    #   lightning
    #   mteb
    #   perceptron
-    #   rich-toolkit
    #   terratorch
    #   typer
-rich-toolkit==0.19.7
-    # via
-    #   fastapi-cli
-    #   fastapi-cloud-cli
-rignore==0.7.6
-    # via fastapi-cloud-cli
 rioxarray==0.22.0
    # via terratorch
 rouge-score==0.1.2
@@ -1286,20 +1070,12 @@ sentence-transformers==5.3.0
    # via
    #   -r requirements/rocm-test.in
    #   mteb
-sentencepiece==0.2.1
-    # via -r requirements/common.txt
 sentry-sdk==2.55.0
-    # via
-    #   fastapi-cloud-cli
-    #   wandb
-setproctitle==1.3.7
-    # via -r requirements/common.txt
+    # via wandb
 setuptools==79.0.1
    # via
    #   -c requirements/common.txt
    #   -c requirements/rocm.txt
-    #   -r requirements/common.txt
-    #   model-hosting-container-standards
    #   pytablewriter
    #   tensorboard
    #   torch
@@ -1316,7 +1092,6 @@ simplejson==3.20.2
 six==1.17.0
    # via
    #   -c requirements/common.txt
-    #   -r requirements/common.txt
    #   junit-xml
    #   lightly
    #   opencensus
@@ -1329,9 +1104,8 @@ smmap==5.0.3
    # via gitdb
 sniffio==1.3.1
    # via
-    #   anthropic
+    #   anyio
    #   httpx
-    #   openai
 sortedcontainers==2.4.0
    # via hypothesis
 soundfile==0.13.1
@@ -1350,16 +1124,10 @@ sqlalchemy==2.0.48
    #   optuna
 sqlitedict==2.1.0
    # via lm-eval
-sse-starlette==3.3.4
-    # via mcp
 starlette==0.52.1
    # via
    #   fastapi
-    #   mcp
-    #   model-hosting-container-standards
-    #   prometheus-fastapi-instrumentator
    #   schemathesis
-    #   sse-starlette
    #   starlette-testclient
 starlette-testclient==0.4.1
    # via schemathesis
@@ -1369,8 +1137,6 @@ stringzilla==4.6.0
    # via albucore
 structlog==25.5.0
    # via gpt-oss
-supervisor==4.3.0
-    # via model-hosting-container-standards
 sympy==1.14.0
    # via
    #   einx
@@ -1414,7 +1180,6 @@ tifffile==2026.3.3
 tiktoken==0.12.0
    # via
    #   -c requirements/common.txt
-    #   -r requirements/common.txt
    #   gpt-oss
    #   lm-eval
    #   mistral-common
@@ -1429,7 +1194,6 @@ timm==1.0.17
 tokenizers==0.22.0
    # via
    #   -c requirements/common.txt
-    #   -r requirements/common.txt
    #   -r requirements/rocm-test.in
    #   transformers
 tomli==2.4.0
@@ -1448,10 +1212,8 @@ torchmetrics==1.9.0
    #   torchgeo
 tqdm==4.67.3
    # via
-    #   -r requirements/common.txt
    #   datasets
    #   evaluate
-    #   gguf
    #   huggingface-hub
    #   lightly
    #   lightning
@@ -1459,7 +1221,6 @@ tqdm==4.67.3
    #   mteb
    #   nltk
    #   open-clip-torch
-    #   openai
    #   optuna
    #   peft
    #   pqdm
@@ -1472,14 +1233,11 @@ tqdm==4.67.3
 transformers==4.57.5
    # via
    #   -c requirements/common.txt
-    #   -r requirements/common.txt
    #   -r requirements/rocm-test.in
-    #   compressed-tensors
    #   genai-perf
    #   peft
    #   sentence-transformers
    #   transformers-stream-generator
-    #   xgrammar
 transformers-stream-generator==0.0.5
    # via -r requirements/rocm-test.in
 tritonclient==2.66.0
@@ -1493,8 +1251,6 @@ typepy==1.3.4
    #   tabledata
 typer==0.24.1
    # via
-    #   fastapi-cli
-    #   fastapi-cloud-cli
    #   fastsafetensors
    #   perceptron
 typeshed-client==2.9.0
@@ -1502,12 +1258,9 @@ typeshed-client==2.9.0
 typing-extensions==4.15.0
    # via
    #   -c requirements/common.txt
-    #   -r requirements/common.txt
    #   aiosignal
    #   albumentations
    #   alembic
-    #   anthropic
-    #   anyio
    #   azure-core
    #   azure-identity
    #   azure-storage-blob
@@ -1519,13 +1272,9 @@ typing-extensions==4.15.0
    #   lightning
    #   lightning-utilities
    #   lm-eval
-    #   mcp
    #   mistral-common
    #   mteb
-    #   openai
    #   opentelemetry-api
-    #   opentelemetry-exporter-otlp-proto-grpc
-    #   opentelemetry-exporter-otlp-proto-http
    #   opentelemetry-sdk
    #   opentelemetry-semantic-conventions
    #   pqdm
@@ -1534,7 +1283,6 @@ typing-extensions==4.15.0
    #   pydantic-extra-types
    #   pytorch-lightning
    #   referencing
-    #   rich-toolkit
    #   sentence-transformers
    #   sqlalchemy
    #   starlette
@@ -1544,13 +1292,10 @@ typing-extensions==4.15.0
    #   typeshed-client
    #   typing-inspection
    #   wandb
-    #   xgrammar
 typing-inspection==0.4.2
    # via
    #   fastapi
-    #   mcp
    #   pydantic
-    #   pydantic-settings
 tzdata==2025.3
    # via arrow
 uri-template==1.3.0
@@ -1566,14 +1311,7 @@ urllib3==2.6.3
    #   sentry-sdk
    #   tritonclient
 uvicorn==0.42.0
-    # via
-    #   fastapi
-    #   fastapi-cli
-    #   fastapi-cloud-cli
-    #   gpt-oss
-    #   mcp
-uvloop==0.22.1
-    # via uvicorn
+    # via gpt-oss
 vector-quantize-pytorch==1.28.0
    # via -r requirements/rocm-test.in
 virtualenv==21.2.0
@@ -1582,16 +1320,10 @@ vocos==0.1.0
    # via -r requirements/rocm-test.in
 wandb==0.25.1
    # via terratorch
-watchfiles==1.1.1
-    # via
-    #   -r requirements/common.txt
-    #   uvicorn
 wcwidth==0.6.0
    # via ftfy
 webcolors==25.10.0
    # via jsonschema
-websockets==16.0
-    # via uvicorn
 werkzeug==3.1.6
    # via
    #   schemathesis
@@ -1602,10 +1334,6 @@ wrapt==2.1.2
    # via smart-open
 xarray==2026.2.0
    # via rioxarray
-xgrammar==0.1.33
-    # via
-    #   -c requirements/common.txt
-    #   -r requirements/common.txt
 xxhash==3.6.0
    # via
    #   datasets
@@ -1060,6 +1060,8 @@ setup(
        ],  # Required for audio processing
        "video": [],  # Kept for backwards compatibility
        "flashinfer": [],  # Kept for backwards compatibility
+        # Optional deps for AMD FP4 quantization support
+        "petit-kernel": ["petit-kernel"],
        # Optional deps for Helion kernel development
        # NOTE: When updating helion version, also update CI files:
        #   - .buildkite/test_areas/kernels.yaml
@@ -0,0 +1,119 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+"""
+Multi-node integration test for MessageQueue TCP fallback.
+
+Verifies that when writer and readers span separate nodes (Docker containers
+with isolated /dev/shm), `create_from_process_group` correctly detects
+cross-node ranks via `in_the_same_node_as()` and falls back to ZMQ TCP
+transport — and that data actually arrives.
+"""
+
+import numpy as np
+import torch.distributed as dist
+
+from vllm.distributed.device_communicators.shm_broadcast import MessageQueue
+from vllm.distributed.parallel_state import in_the_same_node_as
+
+
+def main():
+    dist.init_process_group(backend="gloo")
+
+    rank = dist.get_rank()
+    world_size = dist.get_world_size()
+    assert world_size >= 2, (
+        f"Need at least 2 ranks across nodes, got world_size={world_size}"
+    )
+
+    # Verify that in_the_same_node_as detects cross-node correctly
+    status = in_the_same_node_as(dist.group.WORLD, source_rank=0)
+    local_count = sum(status)
+    print(
+        f"[Rank {rank}] in_the_same_node_as(source=0): {status}  "
+        f"(local={local_count}/{world_size})"
+    )
+    # With 2 Docker containers (1 proc each), rank 0 and rank 1
+    # should be on different nodes.
+    assert local_count < world_size, (
+        f"Expected cross-node ranks but all {world_size} ranks appear local."
+    )
+
+    # Create MessageQueue
+    writer_rank = 0
+    mq = MessageQueue.create_from_process_group(
+        dist.group.WORLD,
+        max_chunk_bytes=1024 * 1024,  # 1 MiB
+        max_chunks=10,
+        writer_rank=writer_rank,
+    )
+
+    # Verify the transport path selection
+    if rank == writer_rank:
+        print(
+            f"[Rank {rank}] Writer: n_local_reader={mq.n_local_reader}, "
+            f"n_remote_reader={mq.n_remote_reader}"
+        )
+        assert mq.n_remote_reader > 0, (
+            "Writer should have at least 1 remote (TCP) reader in a multi-node setup."
+        )
+    else:
+        if status[rank]:
+            assert mq._is_local_reader, (
+                f"Rank {rank} is on the same node as writer but is not a local reader."
+            )
+            print(f"[Rank {rank}] Reader: local (shared memory)")
+        else:
+            assert mq._is_remote_reader, (
+                f"Rank {rank} is on a different node but is not a remote (TCP) reader."
+            )
+            print(f"[Rank {rank}] Reader: remote (TCP)")
+
+    # Test data transfer: simple objects
+    dist.barrier()
+    if rank == writer_rank:
+        mq.enqueue("hello_from_node0")
+    else:
+        msg = mq.dequeue(timeout=10)
+        assert msg == "hello_from_node0"
+    dist.barrier()
+    print(f"[Rank {rank}] Simple object test passed")
+
+    # Test data transfer: numpy arrays
+    np.random.seed(42)
+    arrays = [
+        np.random.randint(0, 100, size=np.random.randint(100, 5000)) for _ in range(100)
+    ]
+
+    dist.barrier()
+    if rank == writer_rank:
+        for arr in arrays:
+            mq.enqueue(arr)
+    else:
+        for i, expected in enumerate(arrays):
+            received = mq.dequeue(timeout=10)
+            assert np.array_equal(expected, received), (
+                f"Array mismatch at index {i}: "
+                f"expected shape {expected.shape}, got shape {received.shape}"
+            )
+    dist.barrier()
+    print(f"[Rank {rank}] Numpy array test passed")
+
+    # Test data transfer: large payload (> max_chunk_bytes)
+    dist.barrier()
+    big_array = np.zeros(200_000, dtype=np.int64)  # ~1.6 MiB > 1 MiB chunk
+    if rank == writer_rank:
+        mq.enqueue(big_array)
+    else:
+        received = mq.dequeue(timeout=10)
+        assert np.array_equal(big_array, received)
+    dist.barrier()
+    print(f"[Rank {rank}] Large payload test passed")
+
+    # Done -- cleanup
+    dist.barrier()
+    print(f"[Rank {rank}] All MessageQueue TCP multi-node tests passed!")
+    dist.destroy_process_group()
+
+
+if __name__ == "__main__":
+    main()
@@ -1,179 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-"""Tests online quantization."""
-
-import pytest
-import torch
-
-from tests.quantization.utils import (
-    _test_online_quant_peak_mem_impl,
-    is_quant_method_supported,
-)
-from vllm.model_executor.layers.linear import UnquantizedLinearMethod
-from vllm.model_executor.layers.quantization.online.fp8 import (
-    Fp8PerBlockOnlineLinearMethod,
-    Fp8PerBlockOnlineMoEMethod,
-    Fp8PerTensorOnlineLinearMethod,
-    Fp8PerTensorOnlineMoEMethod,
-)
-from vllm.platforms import current_platform
-
-
-@pytest.mark.skipif(
-    not is_quant_method_supported("fp8"),
-    reason="FP8 is not supported on this GPU type.",
-)
-@pytest.mark.parametrize(
-    "quant_scheme,online_quant_args,expected_linear_cls,expected_moe_cls",
-    [
-        # simple case - quantization='fp8_per_tensor'
-        (
-            "fp8_per_tensor",
-            None,
-            Fp8PerTensorOnlineLinearMethod,
-            Fp8PerTensorOnlineMoEMethod,
-        ),
-        # simple case - quantization='fp8_per_block'
-        (
-            "fp8_per_block",
-            None,
-            Fp8PerBlockOnlineLinearMethod,
-            Fp8PerBlockOnlineMoEMethod,
-        ),
-        # quantization='online with linear_scheme_override and
-        # moe_scheme_override
-        (
-            "online",
-            {
-                "linear_scheme_override": "fp8_per_block",
-                "moe_scheme_override": "fp8_per_tensor",
-            },
-            Fp8PerBlockOnlineLinearMethod,
-            Fp8PerTensorOnlineMoEMethod,
-        ),
-        # ignore with direct layer name
-        (
-            "fp8_per_tensor",
-            # qkv_proj is fused from q_proj/k_proj/v_proj, so currently the
-            # ignore regex must match the unfused shard names
-            # TODO(future PR): also make 're:.*qkv_proj.*' work
-            {"ignore": ["model.layers.1.self_attn.o_proj", "re:.*[qkv]_proj"]},
-            Fp8PerTensorOnlineLinearMethod,
-            Fp8PerTensorOnlineMoEMethod,
-        ),
-    ],
-)
-@pytest.mark.parametrize(
-    "use_rocm_aiter", [True, False] if current_platform.is_rocm() else [False]
-)
-def test_online_quantization(
-    vllm_runner,
-    quant_scheme: str,
-    online_quant_args: dict | None,
-    expected_linear_cls,
-    expected_moe_cls,
-    use_rocm_aiter: bool,
-    monkeypatch,
-) -> None:
-    """
-    Tests that online quantization frontend configuration works -
-    selecting quant schemes, overriding quant schemes by type, ignoring
-    layers.
-
-    Does not test performance, peak memory usage, etc.
-    """
-
-    if use_rocm_aiter:
-        monkeypatch.setenv("VLLM_ROCM_USE_AITER", "1")
-
-    # `LLM.apply_model` requires pickling a function.
-    monkeypatch.setenv("VLLM_ALLOW_INSECURE_SERIALIZATION", "1")
-
-    # a tiny model with both dense and MoE layers
-    model_name = "ibm-granite/granite-3.0-1b-a400m-base"
-
-    runner_kwargs = dict(
-        quantization=quant_scheme,
-        enforce_eager=True,
-    )
-    if online_quant_args is not None:
-        runner_kwargs["quantization_config"] = online_quant_args
-
-    with vllm_runner(
-        model_name,
-        **runner_kwargs,
-    ) as llm:
-
-        def check_model(model):
-            # checks further down in the test case are hardcoded for this
-            # model
-            assert model_name == "ibm-granite/granite-3.0-1b-a400m-base"
-
-            o_proj = model.model.layers[0].self_attn.o_proj
-            moe = model.model.layers[0].block_sparse_moe.experts
-
-            # o_proj and moe in layer 0 are always quantized (never ignored)
-            # because of how we craft the test case inputs
-            assert isinstance(o_proj.quant_method, expected_linear_cls)
-            if moe is not None:
-                assert isinstance(moe.quant_method, expected_moe_cls)
-
-            if current_platform.is_cuda():
-                assert o_proj.weight.dtype == torch.float8_e4m3fn
-            elif current_platform.is_rocm():
-                assert o_proj.weight.dtype == current_platform.fp8_dtype()
-            else:
-                pytest.skip("Only runs on CUDA and ROCm.")
-
-            # Verify ignored layers are unquantized.
-            if isinstance(online_quant_args, dict) and "ignore" in online_quant_args:
-                # only .*1.self_attn_o_proj is skipped
-                for layer_idx in range(len(model.model.layers)):
-                    o_proj = model.model.layers[layer_idx].self_attn.o_proj
-                    if layer_idx == 1:
-                        assert isinstance(o_proj.quant_method, UnquantizedLinearMethod)
-                    else:
-                        assert isinstance(o_proj.quant_method, expected_linear_cls)
-
-                # every .*self_attn.qkv_proj is skipped
-                for layer_idx in range(len(model.model.layers)):
-                    qkv_proj = model.model.layers[layer_idx].self_attn.qkv_proj
-                    assert isinstance(qkv_proj.quant_method, UnquantizedLinearMethod)
-
-        llm.apply_model(check_model)
-
-        outputs = llm.generate_greedy(["Hello my name is"], max_tokens=4)
-        print(outputs[0][1])
-
-
-@pytest.mark.skipif(
-    not is_quant_method_supported("fp8"),
-    reason="FP8 is not supported on this GPU type.",
-)
-def test_online_quant_peak_mem(
-    vllm_runner,
-    caplog_mp_spawn,
-    monkeypatch,
-) -> None:
-    _test_online_quant_peak_mem_impl(
-        "fp8_per_tensor", vllm_runner, caplog_mp_spawn, monkeypatch
-    )
-
-
-@pytest.mark.skipif(
-    not is_quant_method_supported("fp8"),
-    reason="FP8 is not supported on this GPU type.",
-)
-def test_online_quant_load_format_dummy(
-    vllm_runner,
-    monkeypatch,
-    caplog,
-) -> None:
-    with vllm_runner(
-        "ibm-granite/granite-3.0-1b-a400m-base",
-        quantization="fp8_per_tensor",
-        enforce_eager=True,
-        load_format="dummy",
-    ) as llm:
-        outputs = llm.generate_greedy(["The future of AI is"], max_tokens=4)
-        print(outputs[0][1])
@@ -1,10 +1,6 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project

-import logging
-
-import regex as re
-
 from vllm.model_executor.layers.quantization import get_quantization_config
 from vllm.platforms import current_platform

@@ -25,74 +21,3 @@ def is_quant_method_supported(quant_method: str) -> bool:
    min_capability = get_quantization_config(quant_method).get_min_capability()

    return capability.to_int() >= min_capability
-
-
-def _test_online_quant_peak_mem_impl(
-    quantization_arg_value,
-    vllm_runner,
-    caplog_mp_spawn,
-    monkeypatch,
-) -> None:
-    # Note: `allenai/OLMoE-1B-7B-0125-Instruct` was selected because:
-    # 1. it covers both Linear and MoE paths
-    # 2. it is already used by other tests in CI, so adding it here
-    #    does not increase disk space for CI runners
-    # I really wanted to use `ibm-granite/granite-3.0-1b-a400m-base`
-    # which I think is the smallest MoE model in vLLM (2.5 GiB bf16,
-    # 1.3 GiB fp8), but could not as adding one more model makes CI
-    # run out of disk space.
-    model_name = "allenai/OLMoE-1B-7B-0125-Instruct"
-
-    # Force spawn to ensure caplog_mp_spawn works consistently
-    # (it relies on VLLM_LOGGING_CONFIG_PATH which spawn reads but fork ignores)
-    monkeypatch.setenv("VLLM_WORKER_MULTIPROC_METHOD", "spawn")
-
-    with (
-        caplog_mp_spawn(logging.DEBUG) as log_holder,
-        vllm_runner(
-            model_name,
-            quantization=quantization_arg_value,
-            enforce_eager=True,
-        ) as llm,
-    ):
-        outputs = llm.generate_greedy(["The future of AI is"], max_tokens=4)
-        print(outputs[0][1])
-
-    log_text = log_holder.text
-
-    # Parse memory usage from captured logs
-    model_memory_gib = None
-    peak_memory_gib = None
-    for line in log_text.splitlines():
-        if model_memory_gib is None:
-            match = re.search(r"Model loading took ([\d.]+) GiB memory", line)
-            if match:
-                model_memory_gib = float(match.group(1))
-        if peak_memory_gib is None:
-            match = re.search(
-                r"Peak GPU memory after loading weights: ([\d.]+) GiB", line
-            )
-            if match:
-                peak_memory_gib = float(match.group(1))
-
-    assert model_memory_gib is not None, "Could not find model loading memory log"
-    assert peak_memory_gib is not None, "Could not find peak memory log"
-    print(f"GPU memory used after loading weights: {model_memory_gib} GiB")
-    print(f"Peak GPU memory usage while loading weights: {peak_memory_gib} GiB")
-
-    # model specific, allenai/OLMoE-1B-7B-0125-Instruct fp8 online quant
-    # uses 6.65 GiB for weight loading (bf16 checkpoint is ~12.89 GiB)
-    expected_model_memory_gib = 6.7
-
-    # for allenai/OLMoE-1B-7B-0125-Instruct the number we see today is 9.06
-    # GiB, which is 1.36x above model_memory_gib. A slightly higher number is
-    # expected as when we load and quantize weights in a streaming fashion we
-    # need to have individual weights in bf16 + fp8 alive at the same time.
-    expected_peak_memory_gib = expected_model_memory_gib * 1.4
-
-    assert model_memory_gib < expected_model_memory_gib, (
-        f"{model_memory_gib=} higher than {expected_model_memory_gib}"
-    )
-    assert peak_memory_gib < expected_peak_memory_gib, (
-        f"{peak_memory_gib=} higher than {expected_peak_memory_gib}"
-    )
@@ -10,6 +10,7 @@ from vllm.distributed.kv_transfer.kv_transfer_state import (
    get_kv_transfer_group,
 )
 from vllm.v1.core.sched.output import CachedRequestData, SchedulerOutput
+from vllm.v1.worker.gpu.kv_connector import ActiveKVConnector
 from vllm.v1.worker.kv_connector_model_runner_mixin import KVConnectorModelRunnerMixin

 # Importing utils registers TestExampleConnector with the factory
@@ -59,3 +60,29 @@ def test_kv_connector_mixin_clears_metadata():
    finally:
        # Ensure we clean up the global connector between tests
        ensure_kv_transfer_shutdown()
+
+
+def test_active_kv_connector_runs_lifecycle_hooks_for_empty_metadata():
+    vllm_config = create_vllm_config()
+    vllm_config.kv_transfer_config.kv_connector = "TestExampleConnector"
+    vllm_config.kv_transfer_config.kv_role = "kv_both"
+    vllm_config.kv_transfer_config.kv_connector_extra_config["name"] = "empty"
+
+    ensure_kv_transfer_initialized(vllm_config)
+
+    try:
+        wrapped = get_kv_transfer_group()
+        connector = ActiveKVConnector(vllm_config, {})
+        scheduler_output = _make_empty_scheduler_output()
+
+        connector.pre_forward(scheduler_output)
+        connector.post_forward(scheduler_output)
+
+        assert wrapped.call_record.get("bind_connector_metadata", 0) == 1
+        assert wrapped.call_record.get("handle_preemptions", 0) == 1
+        assert wrapped.call_record.get("start_load_kv", 0) == 1
+        assert wrapped.call_record.get("wait_for_save", 0) == 1
+        assert wrapped.call_record.get("get_finished", 0) == 1
+        assert wrapped.call_record.get("clear_connector_metadata", 0) == 1
+    finally:
+        ensure_kv_transfer_shutdown()
@@ -68,13 +68,10 @@ class IrOpPriorityConfig:
    def set_priority(self):
        """
        Context manager to set the IR op priority for all op members.
-        It also imports IR kernel implementations for the current platform
-        to ensure all implementations are made available.
+        It also imports vllm.kernels to ensure all implementations are made available.
        """
+        import vllm.kernels  # noqa: F401, registers IR op implementations
        from vllm.ir.op import IrOp
-        from vllm.platforms import current_platform
-
-        current_platform.import_ir_kernels()

        with contextlib.ExitStack() as stack:
            for field in fields(self):
@@ -21,7 +21,6 @@ from vllm.config.multimodal import (
    MultiModalConfig,
 )
 from vllm.config.pooler import PoolerConfig
-from vllm.config.quantization import OnlineQuantizationConfigArgs
 from vllm.config.scheduler import RunnerType
 from vllm.config.utils import config, getattr_iter
 from vllm.logger import init_logger
@@ -200,10 +199,6 @@ class ModelConfig:
    `quantization_config` attribute in the model config file. If that is
    `None`, we assume the model weights are not quantized and use `dtype` to
    determine the data type of the weights."""
-    quantization_config: dict[str, Any] | OnlineQuantizationConfigArgs | None = None
-    """Arguments for online quantization.
-    Auto-created when `quantization` equals to one of the string values of
-    the `OnlineQuantScheme` enum."""
    allow_deprecated_quantization: bool = False
    """Whether to allow deprecated quantization methods."""
    enforce_eager: bool = False
@@ -948,6 +943,7 @@ class ModelConfig:
                "modelopt_fp4",
                "modelopt_mxfp8",
                "modelopt_mixed",
+                "petit_nvfp4",
                # Ensure heavy backends are probed last to avoid unnecessary
                # imports during override detection (e.g., MXFP4 imports Triton)
                "mxfp4",
@@ -1,121 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-
-from enum import Enum
-from typing import Any
-
-from pydantic import Field, field_validator
-
-from vllm.config.utils import config
-
-
-class OnlineQuantScheme(Enum):
-    """Supported online quantization schemes."""
-
-    # fp8, weights and activations scaled per-tensor
-    FP8_PER_TENSOR = "fp8_per_tensor"
-
-    # fp8, activations scaled in blocks of 1x128 elements, weights scaled in
-    # blocks of 128x128 elements (popularized by DeepSeek)
-    FP8_PER_BLOCK = "fp8_per_block"
-
-    # TODO(future PRs): add more online quant schemes here: mxfp8, etc
-
-
-@config
-class OnlineQuantizationConfigArgs:
-    """Configuration for online quantization.
-
-    Controls how ``OnlineQuantizationConfig`` is applied to a model.
-    At least one of ``global_scheme``, ``linear_scheme_override``, or
-    ``moe_scheme_override`` must be set.
-    """
-
-    global_scheme: OnlineQuantScheme | None = None
-    """Quantization scheme applied to every supported layer."""
-
-    linear_scheme_override: OnlineQuantScheme | None = None
-    """Quantization scheme override for ``LinearBase`` layers."""
-
-    moe_scheme_override: OnlineQuantScheme | None = None
-    """Quantization scheme override for ``FusedMoE`` layers."""
-
-    ignore: list[str] = Field(default_factory=list)
-    """Layers to skip quantization for. Supports exact names and regex
-    patterns with ``re:`` prefix (e.g. ``re:.*attn.*``), consistent with
-    compressed_tensors layer skipping."""
-
-    @field_validator(
-        "global_scheme", "linear_scheme_override", "moe_scheme_override", mode="before"
-    )
-    @classmethod
-    def _coerce_scheme(
-        cls, v: str | OnlineQuantScheme | None
-    ) -> OnlineQuantScheme | None:
-        if isinstance(v, str):
-            return OnlineQuantScheme(v)
-        return v
-
-
-def resolve_online_quant_config(
-    quantization: str | None,
-    quantization_config: dict[str, Any] | OnlineQuantizationConfigArgs | None,
-) -> OnlineQuantizationConfigArgs | None:
-    """Resolve online quant scheme shorthand into a quantization config.
-
-    If ``quantization`` is an online quant scheme (e.g. ``'fp8_per_tensor'``),
-    ensures ``quantization_config`` has a matching ``global_scheme`` and casts
-    it to :class:`OnlineQuantizationConfigArgs` if needed.
-    """
-    online_quant_values = {s.value for s in OnlineQuantScheme}
-    valid_quantization_values = online_quant_values | {"online"}
-    if quantization not in valid_quantization_values:
-        if quantization_config is not None:
-            raise ValueError(
-                f"quantization_config is only supported when quantization "
-                f"is one of {sorted(valid_quantization_values)}, "
-                f"got quantization={quantization!r}"
-            )
-        return None
-
-    if quantization in online_quant_values:
-        scheme = OnlineQuantScheme(quantization)
-
-        if quantization_config is None:
-            quantization_config = {
-                "global_scheme": scheme.value,
-            }
-        elif isinstance(quantization_config, OnlineQuantizationConfigArgs):
-            if quantization_config.global_scheme is None:
-                quantization_config.global_scheme = scheme
-            elif quantization_config.global_scheme != scheme:
-                raise ValueError(
-                    f"quantization={quantization!r} conflicts with "
-                    f"quantization_config.global_scheme="
-                    f"{quantization_config.global_scheme.value!r}. "
-                    f"These must match when both are specified."
-                )
-        elif isinstance(quantization_config, dict):
-            existing = quantization_config.get("global_scheme")
-            if existing is None:
-                quantization_config["global_scheme"] = scheme.value
-            else:
-                # Coerce to enum for comparison
-                existing_scheme = (
-                    OnlineQuantScheme(existing)
-                    if isinstance(existing, str)
-                    else existing
-                )
-                if existing_scheme != scheme:
-                    raise ValueError(
-                        f"quantization={quantization!r} conflicts "
-                        f"with quantization_config"
-                        f"['global_scheme']={existing!r}. "
-                        f"These must match when both are specified."
-                    )
-
-    # Cast dict to OnlineQuantizationConfigArgs
-    if isinstance(quantization_config, dict):
-        quantization_config = OnlineQuantizationConfigArgs(**quantization_config)
-
-    return quantization_config
@@ -1106,9 +1106,6 @@ class VllmConfig:
            )
        current_platform.check_and_update_config(self)

-        if envs.VLLM_USE_V2_MODEL_RUNNER:
-            self._validate_v2_model_runner()
-
        # Re-compute compile ranges after platform-specific config updates
        # (e.g., XPU may lower max_num_batched_tokens when MLA is enabled)
        self._set_compile_ranges()
@@ -1716,7 +1713,6 @@ class VllmConfig:
            f"dcp_comm_backend={self.parallel_config.dcp_comm_backend}, "  # noqa
            f"disable_custom_all_reduce={self.parallel_config.disable_custom_all_reduce}, "  # noqa
            f"quantization={self.model_config.quantization}, "
-            f"quantization_config={self.model_config.quantization_config}, "  # noqa
            f"enforce_eager={self.model_config.enforce_eager}, "
            f"enable_return_routed_experts={self.model_config.enable_return_routed_experts}, "  # noqa
            f"kv_cache_dtype={self.cache_config.cache_dtype}, "
@@ -1732,49 +1728,6 @@ class VllmConfig:
            f"kernel_config={self.kernel_config!r}"
        )

-    def _validate_v2_model_runner(self) -> None:
-        """Check for features not yet supported by the V2 model runner."""
-        unsupported: list[str] = []
-
-        if self.model_config is not None and self.model_config.has_inner_state:
-            unsupported.append("hybrid/mamba models")
-
-        if self.parallel_config.prefill_context_parallel_size > 1:
-            unsupported.append("prefill context parallelism")
-
-        if (
-            self.speculative_config is not None
-            and self.speculative_config.method not in ("eagle", "eagle3", "mtp")
-        ):
-            unsupported.append(f"speculative method '{self.speculative_config.method}'")
-
-        if self.parallel_config.enable_dbo:
-            unsupported.append("dual batch overlap")
-
-        if (
-            self.model_config is not None
-            and self.model_config.enable_return_routed_experts
-        ):
-            # Will be added by https://github.com/vllm-project/vllm/pull/38163
-            unsupported.append("routed experts capture")
-
-        if self.model_config is not None and self.model_config.logits_processors:
-            unsupported.append("custom logits processors")
-
-        if self.cache_config.kv_sharing_fast_prefill:
-            # Will be added by https://github.com/vllm-project/vllm/pull/35045
-            unsupported.append("KV sharing fast prefill")
-
-        if self.ec_transfer_config is not None:
-            # Will be added by https://github.com/vllm-project/vllm/pull/38390
-            unsupported.append("EC transfer")
-
-        if unsupported:
-            raise ValueError(
-                "VLLM_USE_V2_MODEL_RUNNER does not yet support: "
-                + ", ".join(unsupported)
-            )
-
    def validate_block_size(self) -> None:
        """Validate block_size against DCP and mamba constraints.

@@ -521,7 +521,11 @@ class NixlConnector(KVConnectorBase_V1, SupportsHMA):
    def wait_for_save(self):
        assert self.connector_worker is not None
        assert isinstance(self._connector_metadata, NixlConnectorMetadata)
-        if self.connector_worker.use_host_buffer and self.connector_worker.copy_blocks:
+        if (
+            self.connector_worker.use_host_buffer
+            and self.connector_worker.copy_blocks
+            and self._connector_metadata.reqs_to_save
+        ):
            self.connector_worker.save_kv_to_host(self._connector_metadata)

    def shutdown(self):
@@ -2466,6 +2470,16 @@ class NixlConnectorWorker:
        Start loading by triggering non-blocking nixl_xfer.
        We check for these trnxs to complete in each step().
        """
+        # skip the empty path
+        if (
+            not metadata.reqs_to_recv
+            and not metadata.reqs_to_send
+            and not metadata.reqs_in_batch
+            and not metadata.reqs_not_processed
+            and self._ready_requests.empty()
+        ):
+            return
+
        for req_id, meta in metadata.reqs_to_recv.items():
            meta.local_physical_block_ids = self._logical_to_kernel_block_ids(
                meta.local_block_ids
@@ -112,7 +112,6 @@ from vllm.v1.sample.logits_processor import LogitsProcessor
 from vllm.version import __version__ as VLLM_VERSION

 if TYPE_CHECKING:
-    from vllm.config.quantization import OnlineQuantizationConfigArgs
    from vllm.model_executor.layers.quantization import QuantizationMethods
    from vllm.model_executor.model_loader import LoadFormats
    from vllm.usage.usage_lib import UsageContext
@@ -484,7 +483,6 @@ class EngineArgs:
    hf_overrides: HfOverrides = get_field(ModelConfig, "hf_overrides")
    tokenizer_revision: str | None = ModelConfig.tokenizer_revision
    quantization: QuantizationMethods | str | None = ModelConfig.quantization
-    quantization_config: "dict[str, Any] | OnlineQuantizationConfigArgs | None" = None
    allow_deprecated_quantization: bool = ModelConfig.allow_deprecated_quantization
    enforce_eager: bool = ModelConfig.enforce_eager
    disable_custom_all_reduce: bool = ParallelConfig.disable_custom_all_reduce
@@ -663,12 +661,6 @@ class EngineArgs:
        if isinstance(self.ir_op_priority, dict):
            self.ir_op_priority = IrOpPriorityConfig(**self.ir_op_priority)

-        from vllm.config.quantization import resolve_online_quant_config
-
-        self.quantization_config = resolve_online_quant_config(
-            self.quantization, self.quantization_config
-        )
-
        # Setup plugins
        from vllm.plugins import load_general_plugins

@@ -1439,7 +1431,6 @@ class EngineArgs:
            tokenizer_revision=self.tokenizer_revision,
            max_model_len=self.max_model_len,
            quantization=self.quantization,
-            quantization_config=self.quantization_config,
            allow_deprecated_quantization=self.allow_deprecated_quantization,
            enforce_eager=self.enforce_eager,
            enable_return_routed_experts=self.enable_return_routed_experts,
@@ -34,9 +34,6 @@ from vllm.config.model import (
    RunnerOption,
    TokenizerMode,
 )
-from vllm.config.quantization import (
-    OnlineQuantizationConfigArgs,
-)
 from vllm.distributed.weight_transfer.base import (
    WeightTransferInitRequest,
    WeightTransferUpdateRequest,
@@ -250,9 +247,6 @@ class LLM:
        attention_config: dict[str, Any] | AttentionConfig | None = None,
        kv_cache_memory_bytes: int | None = None,
        compilation_config: int | dict[str, Any] | CompilationConfig | None = None,
-        quantization_config: dict[str, Any]
-        | OnlineQuantizationConfigArgs
-        | None = None,
        logits_processors: list[str | type[LogitsProcessor]] | None = None,
        **kwargs: Any,
    ) -> None:
@@ -373,7 +367,6 @@ class LLM:
            profiler_config=profiler_config_instance,
            attention_config=attention_config_instance,
            compilation_config=compilation_config_instance,
-            quantization_config=quantization_config,
            logits_processors=logits_processors,
            **kwargs,
        )
@@ -154,13 +154,9 @@ is_nvidia_hopper = is_nvidia and (
 )
 use_cuda_graph = is_nvidia and os.environ.get("FLA_USE_CUDA_GRAPH", "0") == "1"
 is_gather_supported = hasattr(triton.language, "gather")
-is_tma_supported = (
-    is_nvidia_hopper
-    and os.getenv("FLA_USE_TMA", "0") == "1"
-    and (
-        hasattr(triton.language, "_experimental_make_tensor_descriptor")
-        or hasattr(triton.language, "make_tensor_descriptor")
-    )
+is_tma_supported = (is_nvidia and torch.cuda.get_device_capability(0)[0] >= 9) and (
+    hasattr(triton.language, "_experimental_make_tensor_descriptor")
+    or hasattr(triton.language, "make_tensor_descriptor")
 )


@@ -79,8 +79,11 @@ class TrtLlmBf16Experts(mk.FusedMoEExpertsMonolithic):
            RoutingMethodType.Default,
            RoutingMethodType.DeepSeekV3,
            RoutingMethodType.Llama4,
-            RoutingMethodType.Renormalize,
-            RoutingMethodType.RenormalizeNaive,
+            # NOTE: TRTLLM Kernel has issue with Qwen3.5 router.
+            # Re-enable once the issue is resolved.
+            # https://github.com/vllm-project/vllm/issues/37591
+            # RoutingMethodType.Renormalize,
+            # RoutingMethodType.RenormalizeNaive
        ]

    @staticmethod
@@ -277,7 +277,13 @@ class TrtLlmFp8ExpertsMonolithic(TrtLlmFp8ExpertsBase, mk.FusedMoEExpertsMonolit
        weight_key: QuantKey | None,
        activation_key: QuantKey | None,
    ) -> bool:
-        """Monolithic kernels need to express router support."""
+        """Monolithic kernels need to express router support.
+        Renormalize/RenormalizeNaive are excluded: the monolithic kernel's
+        internal routing for these methods produces output uncorrelated
+        with the modular kernel's output and with Triton kernel's output
+        for Qwen3.5-35B-A3B-FP8.
+        See: https://github.com/vllm-project/vllm/issues/37591
+        """
        # NOTE(dbari): TopK routing could also be enabled, but need to validate models
        # NOTE(dbari): Default is not implemented and should not be enabled until it is

@@ -289,8 +295,6 @@ class TrtLlmFp8ExpertsMonolithic(TrtLlmFp8ExpertsBase, mk.FusedMoEExpertsMonolit
            return routing_method in [
                RoutingMethodType.DeepSeekV3,
                RoutingMethodType.Simulated,
-                RoutingMethodType.Renormalize,
-                RoutingMethodType.RenormalizeNaive,
            ]
        elif (weight_key, activation_key) == (kFp8StaticTensorSym, kFp8StaticTensorSym):
            # NOTE(dbari): as above, potentially allow others here.
@@ -298,8 +302,6 @@ class TrtLlmFp8ExpertsMonolithic(TrtLlmFp8ExpertsBase, mk.FusedMoEExpertsMonolit
                RoutingMethodType.DeepSeekV3,
                RoutingMethodType.Llama4,
                RoutingMethodType.Simulated,
-                RoutingMethodType.Renormalize,
-                RoutingMethodType.RenormalizeNaive,
            ]
        else:
            raise ValueError("Unsupported quantization scheme.")
@@ -241,12 +241,8 @@ class RMSNorm(CustomOp):
    ) -> torch.Tensor | tuple[torch.Tensor, torch.Tensor]:
        """PyTorch-native implementation equivalent to forward()."""
        if residual is None:
-            # TODO(luka): address the weight=None passing issue more generally
            return ir.ops.rms_norm(
-                x,
-                self.weight.data if self.has_weight else None,
-                self.variance_epsilon,
-                self.variance_size_override,
+                x, self.weight.data, self.variance_epsilon, self.variance_size_override
            )

        return self.forward_static(
@@ -60,6 +60,7 @@ WEIGHT_LOADER_V2_SUPPORTED = [
    "ModelOptFp8PbWoLinearMethod",
    "QuarkLinearMethod",
    "ModelOptNvFp4LinearMethod",
+    "PetitNvFp4LinearMethod",
 ]


@@ -31,14 +31,8 @@ QuantizationMethods = Literal[
    "inc",
    "mxfp4",
    "mxfp8",
+    "petit_nvfp4",
    "cpu_awq",
-    "online",
-    # Below are values of the OnlineQuantScheme enum, specified as strings to
-    # avoid circular import issues. This is here to provide a shortcut where
-    # the user can specify "LLM(..., quantization='fp8_per_tensor')" as
-    # shorthand for creating a more complicated online quant config object
-    "fp8_per_tensor",
-    "fp8_per_block",
 ]
 QUANTIZATION_METHODS: list[str] = list(get_args(QuantizationMethods))

@@ -47,6 +41,7 @@ DEPRECATED_QUANTIZATION_METHODS = [
    "fbgemm_fp8",
    "fp_quant",
    "experts_int8",
+    "petit_nvfp4",
 ]

 # The customized quantization methods which will be added to this dict.
@@ -108,7 +103,6 @@ def get_quantization_config(quantization: str) -> type[QuantizationConfig]:
        raise ValueError(f"Invalid quantization method: {quantization}")

    # lazy import to avoid triggering `torch.compile` too early
-    from vllm.config.quantization import OnlineQuantScheme
    from vllm.model_executor.layers.quantization.quark.quark import QuarkConfig

    from .awq import AWQConfig
@@ -135,7 +129,7 @@ def get_quantization_config(quantization: str) -> type[QuantizationConfig]:
    from .moe_wna16 import MoeWNA16Config
    from .mxfp4 import Mxfp4Config
    from .mxfp8 import Mxfp8Config
-    from .online.base import OnlineQuantizationConfig
+    from .petit import PetitNvFp4Config
    from .torchao import TorchAOConfig

    method_to_config: dict[str, type[QuantizationConfig]] = {
@@ -161,21 +155,9 @@ def get_quantization_config(quantization: str) -> type[QuantizationConfig]:
        "inc": INCConfig,
        "mxfp4": Mxfp4Config,
        "mxfp8": Mxfp8Config,
+        "petit_nvfp4": PetitNvFp4Config,
        "cpu_awq": CPUAWQConfig,
-        "online": OnlineQuantizationConfig,
    }
-
-    # Below are values of the OnlineQuantScheme enum. This is here to provide
-    # a shortcut where the user can specify
-    # "LLM(..., quantization='fp8_per_tensor')" as shorthand for creating a
-    # more complicated online quant config object
-    for scheme in OnlineQuantScheme:
-        assert scheme.value not in method_to_config, (
-            f"Online quant scheme {scheme.value!r} conflicts with an "
-            f"existing quantization method"
-        )
-        method_to_config[scheme.value] = OnlineQuantizationConfig
-
    # Update the `method_to_config` with customized quantization methods.
    method_to_config.update(_CUSTOMIZED_METHOD_TO_QUANT_CONFIG)

@@ -497,8 +497,6 @@ class Fp8LinearMethod(LinearMethodBase):
        return self.fp8_linear.apply_weights(layer, x, bias)


-# TODO(future PR): remove this class in favor of
-# online/fp8.py::Fp8PerTensorOnlineLinearMethod
 class Fp8OnlineLinearMethod(Fp8LinearMethod):
    """Online version of Fp8LinearMethod which loads a full precision checkpoint
    and quantizes weights during loading."""
@@ -921,8 +919,6 @@ class Fp8MoEMethod(FusedMoEMethodBase):
        )


-# TODO(future PR): remove this class in favor of
-# online/fp8.py::Fp8PerTensorOnlineMoEMethod
 class Fp8OnlineMoEMethod(Fp8MoEMethod):
    """MoE method for online FP8 quantization.
    Supports loading quantized FP16/BF16 model checkpoints with dynamic
@@ -1,2 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
@@ -1,116 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-
-from typing import Any
-
-import torch
-
-from vllm.config.quantization import (
-    OnlineQuantizationConfigArgs,
-    OnlineQuantScheme,
-)
-from vllm.model_executor.layers.fused_moe import (
-    FusedMoE,
-)
-from vllm.model_executor.layers.fused_moe.unquantized_fused_moe_method import (
-    UnquantizedFusedMoEMethod,
-)
-from vllm.model_executor.layers.linear import (
-    LinearBase,
-    UnquantizedLinearMethod,
-)
-from vllm.model_executor.layers.quantization import QuantizationMethods
-from vllm.model_executor.layers.quantization.base_config import (
-    QuantizationConfig,
-    QuantizeMethodBase,
-)
-from vllm.model_executor.layers.quantization.compressed_tensors.utils import (
-    should_ignore_layer,
-)
-from vllm.model_executor.layers.quantization.online.fp8 import (
-    Fp8PerBlockOnlineLinearMethod,
-    Fp8PerBlockOnlineMoEMethod,
-    Fp8PerTensorOnlineLinearMethod,
-    Fp8PerTensorOnlineMoEMethod,
-)
-
-
-class OnlineQuantizationConfig(QuantizationConfig):
-    """Model-level config class for online quantization (quantize fp16/bf16 weights
-    during model loading, without requiring a pre-quantized checkpoint)."""
-
-    def __init__(
-        self,
-        args: OnlineQuantizationConfigArgs,
-    ) -> None:
-        super().__init__()
-        if (
-            args.global_scheme is None
-            and args.linear_scheme_override is None
-            and args.moe_scheme_override is None
-        ):
-            raise ValueError(
-                "OnlineQuantizationConfig requires at least one of "
-                "global_scheme, linear_scheme_override, or "
-                "moe_scheme_override to be set."
-            )
-        self.args = args
-        self.quant_scheme = args.global_scheme
-        self.ignored_layers: list[str] = args.ignore
-
-    @classmethod
-    def get_name(cls) -> QuantizationMethods:
-        return "online"
-
-    @classmethod
-    def get_supported_act_dtypes(cls) -> list[torch.dtype]:
-        return [torch.bfloat16, torch.half]
-
-    @classmethod
-    def get_min_capability(cls) -> int:
-        # Note: as more online quant schemes will be added, this
-        # value will become the minimum across all supported schemes.
-        return 75
-
-    @classmethod
-    def get_config_filenames(cls) -> list[str]:
-        return []
-
-    @classmethod
-    def from_config(cls, config: dict[str, Any]) -> "OnlineQuantizationConfig":
-        raise NotImplementedError(
-            "OnlineQuantizationConfig does not support loading from a "
-            "checkpoint config. Use quantization_config or "
-            "quantization='fp8_per_tensor'/'fp8_per_block' instead."
-        )
-
-    def get_quant_method(
-        self, layer: torch.nn.Module, prefix: str
-    ) -> "QuantizeMethodBase | None":
-        if isinstance(layer, LinearBase):
-            if should_ignore_layer(
-                prefix,
-                ignore=self.ignored_layers,
-                fused_mapping=self.packed_modules_mapping,
-            ):
-                return UnquantizedLinearMethod()
-
-            linear_scheme = self.args.linear_scheme_override or self.args.global_scheme
-            if linear_scheme == OnlineQuantScheme.FP8_PER_BLOCK:
-                return Fp8PerBlockOnlineLinearMethod()
-            else:
-                return Fp8PerTensorOnlineLinearMethod()
-        elif isinstance(layer, FusedMoE):
-            if should_ignore_layer(
-                prefix,
-                ignore=self.ignored_layers,
-                fused_mapping=self.packed_modules_mapping,
-            ):
-                return UnquantizedFusedMoEMethod(layer.moe_config)
-
-            moe_scheme = self.args.moe_scheme_override or self.args.global_scheme
-            if moe_scheme == OnlineQuantScheme.FP8_PER_BLOCK:
-                return Fp8PerBlockOnlineMoEMethod(layer=layer)
-            else:
-                return Fp8PerTensorOnlineMoEMethod(layer=layer)
-        return None
@@ -1,632 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-
-from typing import TYPE_CHECKING
-
-import torch
-from torch.nn import Module
-
-if TYPE_CHECKING:
-    import vllm.model_executor.layers.fused_moe.modular_kernel as mk
-    from vllm.model_executor.layers.fused_moe import FusedMoE
-    from vllm.model_executor.layers.fused_moe.config import (
-        FusedMoEConfig,
-        FusedMoEQuantConfig,
-    )
-    from vllm.model_executor.layers.fused_moe.oracle.fp8 import Fp8MoeBackend
-
-import vllm.envs as envs
-from vllm import _custom_ops as ops
-from vllm._aiter_ops import rocm_aiter_ops
-from vllm.model_executor.kernels.linear import init_fp8_linear_kernel
-from vllm.model_executor.layers.fused_moe import (
-    FusedMoEMethodBase,
-)
-from vllm.model_executor.layers.fused_moe.oracle.fp8 import (
-    select_fp8_moe_backend,
-)
-from vllm.model_executor.layers.linear import (
-    LinearMethodBase,
-)
-from vllm.model_executor.layers.quantization.utils.fp8_utils import (
-    W8A8BlockFp8LinearOp,
-    maybe_post_process_fp8_weight_block,
-    process_fp8_weight_block_strategy,
-)
-from vllm.model_executor.layers.quantization.utils.quant_utils import (
-    GroupShape,
-    kFp8Dynamic128Sym,
-    kFp8DynamicTensorSym,
-    kFp8DynamicTokenSym,
-    kFp8Static128BlockSym,
-    kFp8StaticTensorSym,
-)
-from vllm.model_executor.layers.quantization.utils.w8a8_utils import (
-    cutlass_block_fp8_supported,
-    cutlass_fp8_supported,
-)
-from vllm.model_executor.model_loader.reload.layerwise import (
-    initialize_online_processing,
-)
-from vllm.model_executor.parameter import ModelWeightParameter
-from vllm.model_executor.utils import replace_parameter, set_weight_attrs
-from vllm.platforms import current_platform
-from vllm.utils.deep_gemm import is_deep_gemm_supported, per_block_cast_to_fp8
-
-# ---------------------------------------------------------------------------
-# Online FP8 Linear Methods
-# ---------------------------------------------------------------------------
-
-
-class _Fp8OnlineLinearBase(LinearMethodBase):
-    """Shared base for online FP8 linear methods. Loads fp16/bf16 checkpoint
-    weights onto meta device and materializes them just-in-time."""
-
-    uses_meta_device: bool = True
-
-    def create_weights(
-        self,
-        layer: torch.nn.Module,
-        input_size_per_partition: int,
-        output_partition_sizes: list[int],
-        input_size: int,
-        output_size: int,
-        params_dtype: torch.dtype,
-        **extra_weight_attrs,
-    ):
-        output_size_per_partition = sum(output_partition_sizes)
-        weight_loader = extra_weight_attrs.get("weight_loader")
-        layer.logical_widths = output_partition_sizes
-        layer.input_size_per_partition = input_size_per_partition
-        layer.output_size_per_partition = output_size_per_partition
-        layer.orig_dtype = params_dtype
-        layer.weight_block_size = None
-
-        weight = ModelWeightParameter(
-            data=torch.empty(
-                output_size_per_partition,
-                input_size_per_partition,
-                device="meta",  # materialized and processed during loading
-                dtype=params_dtype,
-            ),
-            input_dim=1,
-            output_dim=0,
-            weight_loader=weight_loader,
-        )
-        layer.register_parameter("weight", weight)
-
-        initialize_online_processing(layer)
-
-
-class Fp8PerTensorOnlineLinearMethod(_Fp8OnlineLinearBase):
-    """Online tensorwise FP8 linear quantization.
-    Loads fp16/bf16 weights and quantizes them per-tensor during loading."""
-
-    def __init__(self):
-        self.out_dtype = torch.get_default_dtype()
-
-        # Use per-token quantization for better perf if dynamic and cutlass
-        if cutlass_fp8_supported():
-            activation_quant_key = kFp8DynamicTokenSym
-        else:
-            activation_quant_key = kFp8DynamicTensorSym
-
-        self.fp8_linear = init_fp8_linear_kernel(
-            activation_quant_key=activation_quant_key,
-            weight_quant_key=kFp8StaticTensorSym,
-            out_dtype=torch.get_default_dtype(),
-            module_name=self.__class__.__name__,
-        )
-
-    def process_weights_after_loading(self, layer: Module) -> None:
-        if getattr(layer, "_already_called_process_weights_after_loading", False):
-            return
-
-        layer.input_scale = None
-        qweight, weight_scale = ops.scaled_fp8_quant(layer.weight, scale=None)
-
-        # Update layer with new values.
-        replace_parameter(layer, "weight", qweight.t().data)
-        replace_parameter(layer, "weight_scale", weight_scale.data)
-
-        # Prevent duplicate processing (e.g., during weight reload)
-        layer._already_called_process_weights_after_loading = True
-
-    def apply(
-        self,
-        layer: torch.nn.Module,
-        x: torch.Tensor,
-        bias: torch.Tensor | None = None,
-    ) -> torch.Tensor:
-        # if batch invariant mode is enabled, use BF16 dequant
-        if envs.VLLM_BATCH_INVARIANT:
-            weight_fp8 = layer.weight.to(torch.bfloat16)
-            weight_scale = layer.weight_scale.to(torch.bfloat16)
-            if weight_scale.numel() == 1:
-                # Per-tensor: simple scalar multiplication
-                weight_bf16 = weight_fp8 * weight_scale
-            else:
-                # Multiple scales (fused modules like QKV)
-                if (
-                    weight_scale.dim() == 1
-                    and weight_scale.shape[0] == weight_fp8.shape[0]
-                ):
-                    # Per-row scaling
-                    weight_bf16 = weight_fp8 * weight_scale.unsqueeze(1)
-                else:
-                    # Fallback
-                    weight_bf16 = weight_fp8 * weight_scale
-            return torch.nn.functional.linear(x, weight_bf16.t(), bias)
-
-        return self.fp8_linear.apply_weights(layer, x, bias)
-
-
-class Fp8PerBlockOnlineLinearMethod(_Fp8OnlineLinearBase):
-    """Online blockwise FP8 linear quantization.
-    Loads fp16/bf16 weights and quantizes them per-block during loading."""
-
-    def __init__(self):
-        self.out_dtype = torch.get_default_dtype()
-        self.weight_block_size = [128, 128]
-
-        self.use_deep_gemm = is_deep_gemm_supported()
-        self.use_aiter_and_is_supported = rocm_aiter_ops.is_linear_fp8_enabled()
-        self.cutlass_block_fp8_supported = cutlass_block_fp8_supported()
-
-        self.w8a8_block_fp8_linear = W8A8BlockFp8LinearOp(
-            weight_group_shape=GroupShape(*self.weight_block_size),
-            act_quant_group_shape=GroupShape(1, self.weight_block_size[0]),
-            cutlass_block_fp8_supported=self.cutlass_block_fp8_supported,
-            use_aiter_and_is_supported=self.use_aiter_and_is_supported,
-            use_deep_gemm=self.use_deep_gemm,
-        )
-
-    def create_weights(
-        self,
-        layer: torch.nn.Module,
-        input_size_per_partition: int,
-        output_partition_sizes: list[int],
-        input_size: int,
-        output_size: int,
-        params_dtype: torch.dtype,
-        **extra_weight_attrs,
-    ):
-        super().create_weights(
-            layer,
-            input_size_per_partition,
-            output_partition_sizes,
-            input_size,
-            output_size,
-            params_dtype,
-            **extra_weight_attrs,
-        )
-        layer.weight_block_size = self.weight_block_size
-
-    def process_weights_after_loading(self, layer: Module) -> None:
-        if getattr(layer, "_already_called_process_weights_after_loading", False):
-            return
-
-        layer.input_scale = None
-        block_size = self.weight_block_size
-
-        qweight, weight_scale_inv = per_block_cast_to_fp8(
-            layer.weight, block_size=block_size, use_ue8m0=False
-        )
-
-        qweight, weight_scale_inv = process_fp8_weight_block_strategy(
-            qweight, weight_scale_inv
-        )
-
-        replace_parameter(layer, "weight", qweight.data)
-        replace_parameter(layer, "weight_scale_inv", weight_scale_inv.data)
-
-        maybe_post_process_fp8_weight_block(layer)
-
-        # Prevent duplicate processing (e.g., during weight reload)
-        layer._already_called_process_weights_after_loading = True
-
-    def apply(
-        self,
-        layer: torch.nn.Module,
-        x: torch.Tensor,
-        bias: torch.Tensor | None = None,
-    ) -> torch.Tensor:
-        assert self.weight_block_size is not None
-
-        # Note: batch invariance already handled in the function below
-        return self.w8a8_block_fp8_linear.apply(
-            input=x,
-            weight=layer.weight,
-            weight_scale=layer.weight_scale_inv,
-            input_scale=layer.input_scale,
-            bias=bias,
-        )
-
-
-# ---------------------------------------------------------------------------
-# Online FP8 MoE Methods
-# ---------------------------------------------------------------------------
-
-
-class _Fp8OnlineMoEBase(FusedMoEMethodBase):
-    """Shared base for online FP8 MoE methods. Loads fp16/bf16 checkpoint
-    weights onto meta device and materializes them just-in-time."""
-
-    uses_meta_device: bool = True
-
-    # Declared here for mypy; actual values are set in __init__.
-    fp8_backend: "Fp8MoeBackend"
-    experts_cls: "type[mk.FusedMoEExperts] | None"
-    weight_scale_name: str
-    weight_block_size: list[int] | None
-    moe: "FusedMoEConfig"
-    is_monolithic: bool
-    moe_quant_config: "FusedMoEQuantConfig | None"
-    moe_kernel: "mk.FusedMoEKernel | None"
-
-    def __init__(
-        self,
-        *,
-        weight_block_size: list[int] | None,
-        layer: torch.nn.Module,
-    ):
-        super().__init__(layer.moe_config)
-        self.weight_block_size = weight_block_size
-        self.block_quant: bool = self.weight_block_size is not None
-        self.weight_scale_name = (
-            "weight_scale_inv" if self.block_quant else "weight_scale"
-        )
-
-        # Set weight key and activation key for kernel compatibility
-        if self.block_quant:
-            weight_key = kFp8Static128BlockSym
-            activation_key = kFp8Dynamic128Sym
-        else:
-            weight_key = kFp8StaticTensorSym
-            activation_key = kFp8DynamicTensorSym
-
-        # Select Fp8 MoE backend
-        self.fp8_backend, self.experts_cls = select_fp8_moe_backend(
-            config=self.moe,
-            weight_key=weight_key,
-            activation_key=activation_key,
-            allow_vllm_cutlass=False,
-        )
-
-    def create_weights(
-        self,
-        layer: Module,
-        num_experts: int,
-        hidden_size: int,
-        intermediate_size_per_partition: int,
-        params_dtype: torch.dtype,
-        **extra_weight_attrs,
-    ):
-        layer.num_experts = num_experts
-        layer.orig_dtype = params_dtype
-        layer.weight_block_size = None
-
-        # WEIGHTS
-        w13_weight = torch.nn.Parameter(
-            torch.empty(
-                num_experts,
-                2 * intermediate_size_per_partition,
-                hidden_size,
-                device="meta",
-                dtype=params_dtype,
-            ),
-            requires_grad=False,
-        )
-        layer.register_parameter("w13_weight", w13_weight)
-        set_weight_attrs(w13_weight, extra_weight_attrs)
-
-        w2_weight = torch.nn.Parameter(
-            torch.empty(
-                num_experts,
-                hidden_size,
-                intermediate_size_per_partition,
-                device="meta",  # materialized and processed during loading
-                dtype=params_dtype,
-            ),
-            requires_grad=False,
-        )
-        layer.register_parameter("w2_weight", w2_weight)
-        set_weight_attrs(w2_weight, extra_weight_attrs)
-
-        # BIASES (for models like GPT-OSS that have biased MoE)
-        if self.moe.has_bias:
-            w13_bias = torch.nn.Parameter(
-                torch.zeros(
-                    num_experts,
-                    2 * intermediate_size_per_partition,
-                    device="meta",  # materialized and processed during loading
-                    dtype=layer.orig_dtype,
-                ),
-                requires_grad=False,
-            )
-            layer.register_parameter("w13_bias", w13_bias)
-            set_weight_attrs(w13_bias, extra_weight_attrs)
-
-            w2_bias = torch.nn.Parameter(
-                torch.zeros(
-                    num_experts,
-                    hidden_size,
-                    device="meta",  # materialized and processed during loading
-                    dtype=layer.orig_dtype,
-                ),
-                requires_grad=False,
-            )
-            layer.register_parameter("w2_bias", w2_bias)
-            set_weight_attrs(w2_bias, extra_weight_attrs)
-
-        layer.w13_input_scale = None
-        layer.w2_input_scale = None
-
-        initialize_online_processing(layer)
-
-    def _setup_kernel(
-        self,
-        layer: "FusedMoE",
-        w13: torch.Tensor,
-        w2: torch.Tensor,
-        w13_scale: torch.Tensor,
-        w2_scale: torch.Tensor,
-        w13_input_scale: torch.Tensor | None,
-        w2_input_scale: torch.Tensor | None,
-    ) -> None:
-        from vllm.model_executor.layers.fused_moe.oracle.fp8 import (
-            convert_to_fp8_moe_kernel_format,
-            make_fp8_moe_kernel,
-        )
-
-        # Shuffle weights to runtime format.
-        w13, w2, w13_scale, w2_scale = convert_to_fp8_moe_kernel_format(
-            fp8_backend=self.fp8_backend,
-            layer=layer,
-            w13=w13,
-            w2=w2,
-            w13_scale=w13_scale,
-            w2_scale=w2_scale,
-            w13_input_scale=w13_input_scale,
-            w2_input_scale=w2_input_scale,
-        )
-
-        # Replace parameters with updated versions. Note that this helper
-        # function ensures the replacement is compatible with RL weight reloads.
-        replace_parameter(layer, "w13_weight", w13)
-        replace_parameter(layer, "w2_weight", w2)
-        replace_parameter(layer, f"w13_{self.weight_scale_name}", w13_scale)
-        replace_parameter(layer, f"w2_{self.weight_scale_name}", w2_scale)
-
-        self.moe_quant_config = self.get_fused_moe_quant_config(layer)
-        if self.moe_quant_config:
-            assert self.experts_cls is not None
-            self.moe_kernel = make_fp8_moe_kernel(
-                moe_quant_config=self.moe_quant_config,
-                moe_config=self.moe,
-                fp8_backend=self.fp8_backend,
-                experts_cls=self.experts_cls,
-                routing_tables=layer._maybe_init_expert_routing_tables(),
-                shared_experts=layer.shared_experts,
-            )
-
-    def maybe_make_prepare_finalize(
-        self,
-        routing_tables: tuple[torch.Tensor, torch.Tensor, torch.Tensor] | None = None,
-    ) -> "mk.FusedMoEPrepareAndFinalizeModular | None":
-        raise ValueError(
-            f"{self.__class__.__name__} uses the new modular kernel "
-            "initialization logic. This function should not be called."
-        )
-
-    def get_fused_moe_quant_config(
-        self, layer: torch.nn.Module
-    ) -> "FusedMoEQuantConfig":
-        from vllm.model_executor.layers.fused_moe.oracle.fp8 import (
-            make_fp8_moe_quant_config,
-        )
-
-        w1_scale = getattr(layer, f"w13_{self.weight_scale_name}")
-        w2_scale = getattr(layer, f"w2_{self.weight_scale_name}")
-        a1_scale = layer.w13_input_scale
-        a2_scale = layer.w2_input_scale
-
-        quant_config = make_fp8_moe_quant_config(
-            fp8_backend=self.fp8_backend,
-            w1_scale=w1_scale,
-            w2_scale=w2_scale,
-            a1_scale=a1_scale,
-            a2_scale=a2_scale,
-            block_shape=self.weight_block_size,
-        )
-
-        # Inject biases into the quant config if the model has them
-        # (e.g. GPT-OSS biased MoE)
-        if quant_config is not None and self.moe.has_bias:
-            w13_bias = getattr(layer, "w13_bias", None)
-            w2_bias = getattr(layer, "w2_bias", None)
-            if w13_bias is not None:
-                quant_config._w1.bias = w13_bias
-            if w2_bias is not None:
-                quant_config._w2.bias = w2_bias
-
-        return quant_config
-
-    @property
-    def supports_eplb(self) -> bool:
-        return True
-
-    def apply_monolithic(
-        self,
-        layer: "FusedMoE",
-        x: torch.Tensor,
-        router_logits: torch.Tensor,
-    ) -> torch.Tensor | tuple[torch.Tensor, torch.Tensor]:
-        assert self.is_monolithic
-        assert self.moe_kernel is not None
-        return self.moe_kernel.apply_monolithic(
-            x,
-            layer.w13_weight,
-            layer.w2_weight,
-            router_logits,
-            activation=layer.activation,
-            global_num_experts=layer.global_num_experts,
-            expert_map=layer.expert_map,
-            apply_router_weight_on_input=layer.apply_router_weight_on_input,
-            num_expert_group=layer.num_expert_group,
-            topk_group=layer.topk_group,
-            e_score_correction_bias=layer.e_score_correction_bias,
-            routed_scaling_factor=layer.routed_scaling_factor,
-        )
-
-    def apply(
-        self,
-        layer: "FusedMoE",
-        x: torch.Tensor,
-        topk_weights: torch.Tensor,
-        topk_ids: torch.Tensor,
-        shared_experts_input: torch.Tensor | None,
-    ) -> torch.Tensor | tuple[torch.Tensor, torch.Tensor]:
-        assert not self.is_monolithic
-        assert self.moe_kernel is not None
-        return self.moe_kernel.apply(
-            x,
-            layer.w13_weight,
-            layer.w2_weight,
-            topk_weights,
-            topk_ids,
-            activation=layer.activation,
-            global_num_experts=layer.global_num_experts,
-            expert_map=layer.expert_map,
-            apply_router_weight_on_input=layer.apply_router_weight_on_input,
-            shared_experts_input=shared_experts_input,
-        )
-
-
-class Fp8PerTensorOnlineMoEMethod(_Fp8OnlineMoEBase):
-    """Online tensorwise FP8 MoE quantization.
-    Loads fp16/bf16 weights and quantizes them per-tensor during loading."""
-
-    def __init__(
-        self,
-        *,
-        layer: torch.nn.Module,
-    ):
-        super().__init__(
-            weight_block_size=None,
-            layer=layer,
-        )
-
-    def process_weights_after_loading(self, layer: Module) -> None:
-        # TODO(@ksayers): inplace fp8 quant kernel, initialize scales with ones
-        if getattr(layer, "_already_called_process_weights_after_loading", False):
-            return
-
-        # If checkpoint is fp16, quantize in place.
-        fp8_dtype = current_platform.fp8_dtype()
-        w13 = torch.empty_like(layer.w13_weight, dtype=fp8_dtype)
-        w2 = torch.empty_like(layer.w2_weight, dtype=fp8_dtype)
-        w13_scale = torch.ones(
-            layer.num_experts, device=w13.device, dtype=torch.float32
-        )
-        w2_scale = torch.ones(layer.num_experts, device=w2.device, dtype=torch.float32)
-        layer.w13_input_scale = None
-        layer.w2_input_scale = None
-
-        for expert in range(layer.local_num_experts):
-            w13[expert, :, :], w13_scale[expert] = ops.scaled_fp8_quant(
-                layer.w13_weight[expert, :, :]
-            )
-            w2[expert, :, :], w2_scale[expert] = ops.scaled_fp8_quant(
-                layer.w2_weight[expert, :, :]
-            )
-
-        # Shuffle weights to runtime format and setup kernel.
-        self._setup_kernel(
-            layer,
-            w13,
-            w2,
-            w13_scale,
-            w2_scale,
-            w13_input_scale=layer.w13_input_scale,
-            w2_input_scale=layer.w2_input_scale,
-        )
-
-        # Prevent duplicate processing (e.g., during weight reload)
-        layer._already_called_process_weights_after_loading = True
-
-
-class Fp8PerBlockOnlineMoEMethod(_Fp8OnlineMoEBase):
-    """Online blockwise FP8 MoE quantization.
-    Loads fp16/bf16 weights and quantizes them per-block during loading."""
-
-    def __init__(
-        self,
-        *,
-        layer: torch.nn.Module,
-    ):
-        super().__init__(
-            weight_block_size=[128, 128],
-            layer=layer,
-        )
-
-    def process_weights_after_loading(self, layer: Module) -> None:
-        if getattr(layer, "_already_called_process_weights_after_loading", False):
-            return
-
-        fp8_dtype = current_platform.fp8_dtype()
-        w13 = torch.empty_like(layer.w13_weight, dtype=fp8_dtype)
-        w2 = torch.empty_like(layer.w2_weight, dtype=fp8_dtype)
-
-        block_size = self.weight_block_size
-        assert block_size is not None
-        block_n, block_k = block_size
-
-        # Create block-shaped scales (computed here rather than in
-        # create_weights because online quant doesn't need them until now).
-        num_experts = layer.local_num_experts
-        _, w13_out, w13_in = layer.w13_weight.shape
-        _, w2_out, w2_in = layer.w2_weight.shape
-
-        w13_scale = torch.ones(
-            num_experts,
-            (w13_out + block_n - 1) // block_n,
-            (w13_in + block_k - 1) // block_k,
-            dtype=torch.float32,
-            device=w13.device,
-        )
-        w2_scale = torch.ones(
-            num_experts,
-            (w2_out + block_n - 1) // block_n,
-            (w2_in + block_k - 1) // block_k,
-            dtype=torch.float32,
-            device=w2.device,
-        )
-
-        for expert in range(num_experts):
-            w13[expert], w13_scale[expert] = per_block_cast_to_fp8(
-                layer.w13_weight[expert],
-                block_size=block_size,
-                use_ue8m0=False,
-            )
-            w2[expert], w2_scale[expert] = per_block_cast_to_fp8(
-                layer.w2_weight[expert],
-                block_size=block_size,
-                use_ue8m0=False,
-            )
-
-        layer.weight_block_size = block_size
-
-        # Shuffle weights to runtime format and setup kernel.
-        self._setup_kernel(
-            layer,
-            w13,
-            w2,
-            w13_scale,
-            w2_scale,
-            layer.w13_input_scale,
-            layer.w2_input_scale,
-        )
-
-        # Prevent duplicate processing (e.g., during weight reload)
-        layer._already_called_process_weights_after_loading = True
@@ -0,0 +1,319 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+# Adapted from https://github.com/vllm-project/vllm/blob/main/vllm/model_executor/layers/quantization/modelopt.py
+
+from typing import Any
+
+import regex as re
+import torch
+from torch.nn.parameter import Parameter
+
+from vllm.logger import init_logger
+from vllm.model_executor.layers.attention import Attention
+from vllm.model_executor.layers.linear import (
+    LinearBase,
+    LinearMethodBase,
+    UnquantizedLinearMethod,
+)
+from vllm.model_executor.layers.quantization import QuantizationMethods
+from vllm.model_executor.layers.quantization.base_config import (
+    QuantizationConfig,
+    QuantizeMethodBase,
+)
+from vllm.model_executor.layers.quantization.kv_cache import BaseKVCacheMethod
+from vllm.model_executor.layers.quantization.utils.petit_utils import (
+    apply_petit_nvfp4_linear,
+    prepare_nvfp4_layer_for_petit,
+    verify_petit_nvfp4_supported,
+)
+from vllm.model_executor.layers.quantization.utils.quant_utils import is_layer_skipped
+from vllm.model_executor.parameter import ModelWeightParameter, PerTensorScaleParameter
+from vllm.platforms import current_platform
+
+# Initialize logger for the module
+logger = init_logger(__name__)
+
+
+# Configuration class to support the NVFP4 quantized model
+# generated by the ModelOpt quantization tool
+class PetitNvFp4Config(QuantizationConfig):
+    """Config class for Petit FP4."""
+
+    def __init__(
+        self,
+        is_checkpoint_nvfp4_serialized: bool = False,
+        kv_cache_quant_algo: str | None = None,
+        group_size: int | None = None,
+        exclude_modules: list[str] | None = None,
+    ) -> None:
+        self._check_hardware_support()
+        self.is_checkpoint_nvfp4_serialized = is_checkpoint_nvfp4_serialized
+        if is_checkpoint_nvfp4_serialized:
+            logger.warning(
+                "Detected nvfp4 checkpoint. Please note that the "
+                "format is experimental and subject to change."
+            )
+        self.group_size = group_size
+        self.kv_cache_quant_algo = kv_cache_quant_algo
+        self.exclude_modules = exclude_modules
+
+    def _check_hardware_support(self) -> None:
+        """
+        Verifies that the current hardware is supported by the Petit backend.
+        This backend is specifically designed for AMD GPUs and is not
+        supported on the CUDA platform.
+        """
+        # This check ensures the code is NOT running on an NVIDIA GPU.
+        if current_platform.is_cuda():
+            raise ValueError(
+                "The 'petit' quantization backend is designed for AMD GPUs "
+                "and is not supported on the CUDA platform. For NVIDIA GPUs, "
+                "please use a different quantization method such as FP8, AWQ, "
+                "or GPTQ."
+            )
+
+    @classmethod
+    def get_name(cls) -> QuantizationMethods:
+        return "petit_nvfp4"
+
+    @classmethod
+    def get_supported_act_dtypes(cls) -> list[torch.dtype]:
+        return [torch.bfloat16, torch.half]
+
+    @classmethod
+    def get_min_capability(cls) -> int:
+        # Petit supports the gfx90a and gfx942 GPUs
+        return 90
+
+    @classmethod
+    def get_config_filenames(cls) -> list[str]:
+        return ["hf_quant_config.json"]
+
+    @classmethod
+    def from_config(cls, config: dict[str, Any]) -> "PetitNvFp4Config":
+        qc = cls.get_from_keys(config, ["quantization"])
+
+        quant_method_raw = qc.get("quant_algo")
+        if not isinstance(quant_method_raw, str) or not quant_method_raw:
+            raise ValueError("Missing or invalid 'quant_algo' in quantization config.")
+        quant_method = quant_method_raw.upper()
+
+        group_size_raw = qc.get("group_size")
+        if not isinstance(group_size_raw, int):
+            raise ValueError(
+                "Missing or invalid 'group_size' (int) in hf_quant_config.json."
+            )
+        group_size = group_size_raw
+
+        verify_petit_nvfp4_supported(quant_method, group_size)
+
+        kv_cache_quant_algo_raw = qc.get("kv_cache_quant_algo") or "auto"
+        if not isinstance(kv_cache_quant_algo_raw, str):
+            raise ValueError("'kv_cache_quant_algo' must be a string if provided.")
+        kv_cache_quant_algo = kv_cache_quant_algo_raw
+
+        exclude_raw = qc.get("exclude_modules", [])
+        if exclude_raw is None:
+            exclude_modules: list[str] = []
+        elif isinstance(exclude_raw, list) and all(
+            isinstance(x, str) for x in exclude_raw
+        ):
+            exclude_modules = exclude_raw
+        else:
+            raise ValueError("'exclude_modules' must be a list[str] (or omitted).")
+
+        is_checkpoint_nvfp4_serialized = "NVFP4" in quant_method
+
+        return cls(
+            is_checkpoint_nvfp4_serialized=is_checkpoint_nvfp4_serialized,
+            kv_cache_quant_algo=kv_cache_quant_algo,
+            group_size=group_size,
+            exclude_modules=exclude_modules,
+        )
+
+    @classmethod
+    def override_quantization_method(
+        cls, hf_quant_cfg, user_quant
+    ) -> QuantizationMethods | None:
+        if not current_platform.is_rocm():
+            return None
+
+        qc = hf_quant_cfg.get("quantization", hf_quant_cfg)
+        algo = (qc.get("quant_algo") or qc.get("quant_method") or "").upper()
+        if algo in ("NVFP4", "MODELOPT_FP4", "MODELOPT"):
+            return cls.get_name()  # "petit_nvfp4"
+        return None
+
+    @classmethod
+    def is_petit_nvfp4_compatible(cls, quant_config: dict[str, Any]) -> bool:
+        qc = quant_config.get("quantization", quant_config)
+        algo = (qc.get("quant_algo") or qc.get("quant_method") or "").upper()
+        return algo == "NVFP4"
+
+    def is_layer_excluded(self, prefix: str, exclude_modules: list[str]) -> bool:
+        for pattern in exclude_modules:
+            regex_str = pattern.replace(".", r"\.").replace("*", r".*")
+            if re.fullmatch(regex_str, prefix):
+                return True
+        return False
+
+    def get_quant_method(
+        self, layer: torch.nn.Module, prefix: str
+    ) -> "QuantizeMethodBase | None":
+        exclude = self.require_exclude_modules()
+
+        if isinstance(layer, LinearBase):
+            if is_layer_skipped(prefix, exclude) or self.is_layer_excluded(
+                prefix, exclude
+            ):
+                return UnquantizedLinearMethod()
+            return PetitNvFp4LinearMethod(self)
+        elif isinstance(layer, Attention):
+            return PetitFp8KVCacheMethod(self)
+        return None
+
+    def get_scaled_act_names(self) -> list[str]:
+        return []
+
+    def require_group_size(self) -> int:
+        if self.group_size is None:
+            logger.warning("group_size not set; defaulting to 16 for NVFP4.")
+            return 16
+        return self.group_size
+
+    def require_kv_cache_quant_algo(self) -> str:
+        return self.kv_cache_quant_algo or "auto"
+
+    def require_exclude_modules(self) -> list[str]:
+        return list(self.exclude_modules or [])
+
+
+class PetitFp8KVCacheMethod(BaseKVCacheMethod):
+    """
+    Supports loading kv-cache scaling factors from FP8 checkpoints.
+    """
+
+    def __init__(self, quant_config: PetitNvFp4Config):
+        super().__init__(quant_config)
+
+
+class PetitNvFp4LinearMethod(LinearMethodBase):
+    """Linear method for NVFP4.
+    Supports loading NVFP4 checkpoints with the following structure:
+
+    |Tensor Name           | datatype      |  shape      |
+    |----------------------------------------------------|
+    |input_scale           | torch.float32 | scalar      |
+    |weight                | NVFP4(SE2M1)  | [1, X, y/2] |
+    |weight_scale          | FP8-E4M3      | [X, Y]      |
+    |weight_scale_2        | torch.float32 | scalar      |
+
+    The weights are quantized per block of 16 elements.
+    Args: quant_config: The ModelOpt quantization config.
+    """
+
+    def __init__(self, quant_config: PetitNvFp4Config):
+        self.quant_config = quant_config
+
+    def create_weights(
+        self,
+        layer: torch.nn.Module,
+        input_size_per_partition: int,
+        output_partition_sizes: list[int],
+        input_size: int,
+        output_size: int,
+        params_dtype: torch.dtype,
+        **extra_weight_attrs,
+    ):
+        del input_size, output_size
+        if not self.quant_config.is_checkpoint_nvfp4_serialized:
+            raise ValueError(
+                "NVFP4 quantization was selected, "
+                " dynamic quantization is not supported."
+            )
+
+        output_size_per_partition = sum(output_partition_sizes)
+        weight_loader = extra_weight_attrs.get("weight_loader")
+
+        layer.logical_widths = output_partition_sizes
+
+        layer.input_size_per_partition = input_size_per_partition
+        layer.output_size_per_partition = output_size_per_partition
+        if input_size_per_partition % 16 != 0:
+            raise ValueError(
+                "Unsupported model when in features size is not multiple of 16"
+            )
+
+        weight_dtype = (
+            torch.float8_e4m3fn
+            if self.quant_config.is_checkpoint_nvfp4_serialized
+            else params_dtype
+        )
+
+        weight = ModelWeightParameter(
+            data=torch.empty(
+                # 2 fp4 data is packed in one uint8 in the input dimension
+                output_size_per_partition,
+                input_size_per_partition // 2,
+                dtype=torch.uint8,
+            ),
+            input_dim=1,
+            output_dim=0,
+            weight_loader=weight_loader,
+        )
+        layer.register_parameter("weight", weight)
+
+        input_scale = PerTensorScaleParameter(
+            data=torch.empty(len(output_partition_sizes), dtype=torch.float32),
+            weight_loader=weight_loader,
+        )
+
+        layer.register_parameter("input_scale", input_scale)
+
+        weight_scale_2 = PerTensorScaleParameter(
+            data=torch.empty(len(output_partition_sizes), dtype=torch.float32),
+            weight_loader=weight_loader,
+        )
+        layer.register_parameter("weight_scale_2", weight_scale_2)
+
+        group_size = self.quant_config.require_group_size()
+        weight_scale = ModelWeightParameter(
+            data=torch.empty(
+                output_size_per_partition,
+                input_size_per_partition // group_size,
+                dtype=weight_dtype,
+            ),
+            input_dim=1,
+            output_dim=0,
+            weight_loader=weight_loader,
+        )
+
+        layer.register_parameter("weight_scale", weight_scale)
+
+    def process_weights_after_loading(self, layer: torch.nn.Module) -> None:
+        input_scale_2 = layer.input_scale.max().to(torch.float32)
+        weight_scale_2 = layer.weight_scale_2.max().to(torch.float32)
+        layer.input_scale = Parameter(input_scale_2, requires_grad=False)
+        layer.weight_scale_2 = Parameter(weight_scale_2, requires_grad=False)
+        layer.alpha = Parameter(
+            layer.input_scale * layer.weight_scale_2, requires_grad=False
+        )
+
+        prepare_nvfp4_layer_for_petit(layer)
+        del layer.input_scale
+
+    def apply(
+        self,
+        layer: torch.nn.Module,
+        x: torch.Tensor,
+        bias: torch.Tensor | None = None,
+    ) -> torch.Tensor:
+        return apply_petit_nvfp4_linear(
+            input=x,
+            weight=layer.weight,
+            weight_scale=layer.weight_scale,
+            weight_scale_2=layer.weight_scale_2,
+            size_n=layer.output_size_per_partition,
+            size_k=layer.input_size_per_partition,
+            bias=bias,
+        )
@@ -0,0 +1,120 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+from typing import TYPE_CHECKING
+
+import torch
+
+# TYPE_CHECKING is used for static type analysis to prevent circular imports.
+if TYPE_CHECKING:
+    from types import ModuleType
+
+# 1. Create a global variable as a placeholder for the module
+_petit_kernel: "ModuleType | None" = None
+
+_PETIT_INSTALL_MSG = (
+    "Petit is not installed. Please install it with `pip install petit-kernel`."
+)
+
+
+def _import_petit_kernel() -> "ModuleType":
+    """
+    A helper function to handle the lazy import.
+    The first time this function is called, it will import the petit_kernel
+    library and store it in the global _petit_kernel variable.
+    Subsequent calls will return the already-loaded module directly.
+    """
+    global _petit_kernel
+    if _petit_kernel is not None:
+        return _petit_kernel
+
+    try:
+        import petit_kernel
+
+        _petit_kernel = petit_kernel
+        return _petit_kernel
+    except ImportError:
+        # The 'from None' syntax prevents chaining the original ImportError,
+        # making the traceback cleaner.
+        raise ImportError(_PETIT_INSTALL_MSG) from None
+
+
+def _check_petit_nvfp4_supported(
+    quant_method: str, group_size: int | None
+) -> tuple[bool, str | None]:
+    if quant_method != "NVFP4":
+        return (
+            False,
+            (
+                "Petit currently only supports: NVFP4 quantizations in sglang. "
+                "Please check the `hf_quant_config.json` file for your model's "
+                "quant configuration."
+            ),
+        )
+    if group_size is not None and group_size != 16:
+        return (
+            False,
+            "Petit currently only supports: group_size=16 quantizations.",
+        )
+    return (True, None)
+
+
+def verify_petit_nvfp4_supported(quant_method: str, group_size: int | None) -> None:
+    supported, error_msg = _check_petit_nvfp4_supported(quant_method, group_size)
+    if not supported:
+        assert error_msg is not None
+        raise ValueError(error_msg)
+
+
+def prepare_nvfp4_layer_for_petit(layer: torch.nn.Module) -> None:
+    # 2. Call _import_petit_kernel() to trigger (or get) the import.
+    petit_kernel = _import_petit_kernel()
+
+    # Repack weights to petit format
+    part_size_n = layer.output_size_per_partition
+    part_size_k = layer.input_size_per_partition
+    qweight = layer.weight.view(torch.int32).contiguous()
+
+    # 3. Call functions through the imported module variable.
+    petit_qweight = petit_kernel.repack_nvfp4(
+        qweight, size_n=part_size_n, size_k=part_size_k
+    )
+    layer.weight = torch.nn.Parameter(petit_qweight, requires_grad=False)
+
+    # Permute scales
+    weight_scale = petit_kernel.process_nvfp4_scales(
+        scales=layer.weight_scale, size_k=part_size_k, size_n=part_size_n
+    )
+    layer.weight_scale = torch.nn.Parameter(weight_scale, requires_grad=False)
+
+
+def apply_petit_nvfp4_linear(
+    input: torch.Tensor,
+    weight: torch.Tensor,
+    weight_scale: torch.Tensor,
+    weight_scale_2: torch.Tensor,
+    size_n: int,
+    size_k: int,
+    bias: torch.Tensor | None = None,
+) -> torch.Tensor:
+    # Trigger (or get) the import here as well.
+    petit_kernel = _import_petit_kernel()
+
+    reshaped_x = input.reshape(-1, input.shape[-1])
+    out_shape = input.shape[:-1] + (size_n,)
+
+    # TODO: Use auto-tuning to find the performant solution_id
+    # Call the function via the module variable.
+    output = petit_kernel.mul_nvfp4_a16(
+        a=reshaped_x,
+        b=weight,
+        s=weight_scale,
+        global_scale=weight_scale_2,
+        size_m=reshaped_x.size(0),
+        size_n=size_n,
+        size_k=size_k,
+        solution_id=-1,
+    )
+    if bias is not None:
+        output.add_(bias)  # In-place add
+
+    return output.reshape(out_shape)
@@ -296,13 +296,6 @@ def get_quant_config(
        )

    if hf_quant_config is not None:
-        if model_config.quantization_config is not None:
-            raise ValueError(
-                "Setting `quantization_config` for online "
-                "quantization when the model checkpoint already "
-                "has a `quantization_config` is not supported"
-            )
-
        # For modelopt_mixed, config.json's quantization_config may or may
        # not contain the per-layer quantized_layers map.  Newer checkpoints
        # embed it directly; older ones keep it only in hf_quant_config.json.
@@ -326,12 +319,6 @@ def get_quant_config(
    quantization_config_file = hf_overrides.get("quantization_config_file", None)
    if quantization_config_file is not None:
        if hasattr(quant_cls, "from_config_file"):
-            if model_config.quantization_config is not None:
-                raise ValueError(
-                    "Setting `quantization_config` for online "
-                    "quantization when the model checkpoint already "
-                    "has a `quantization_config` is not supported"
-                )
            return quant_cls.from_config_file(quantization_config_file)
        else:
            raise NotImplementedError(
@@ -342,12 +329,6 @@ def get_quant_config(
    quantization_config_json = hf_overrides.get("quantization_config_dict_json", None)
    if quantization_config_json is not None:
        if hasattr(quant_cls, "from_config_dict_json"):
-            if model_config.quantization_config is not None:
-                raise ValueError(
-                    "Setting `quantization_config` for online "
-                    "quantization when the model checkpoint already "
-                    "has a `quantization_config` is not supported"
-                )
            return quant_cls.from_config_dict_json(quantization_config_json)
        else:
            raise NotImplementedError(
@@ -356,19 +337,6 @@ def get_quant_config(
                f"{quant_cls}"
            )

-    # Online quantization doesn't read from checkpoint configs — it quantizes
-    # fp16/bf16 weights on the fly during loading.
-    if model_config.quantization_config is not None:
-        from vllm.config.quantization import OnlineQuantizationConfigArgs
-        from vllm.model_executor.layers.quantization.online.base import (
-            OnlineQuantizationConfig,
-        )
-
-        assert isinstance(
-            model_config.quantization_config, OnlineQuantizationConfigArgs
-        )
-        return OnlineQuantizationConfig(args=model_config.quantization_config)
-
    # Inflight BNB quantization
    if model_config.quantization == "bitsandbytes":
        return quant_cls.from_config({})
@@ -16,6 +16,7 @@ from vllm.distributed import (
    get_tensor_model_parallel_world_size,
    tensor_model_parallel_all_reduce,
 )
+from vllm.logger import init_logger
 from vllm.model_executor.layers.activation import SiluAndMul
 from vllm.model_executor.layers.attention import Attention
 from vllm.model_executor.layers.fused_moe import fused_experts, fused_topk
@@ -41,7 +42,6 @@ from vllm.transformers_utils.configs.arctic import ArcticConfig

 from .interfaces import SupportsPP, SupportsQuant
 from .utils import (
-    AutoWeightsLoader,
    extract_layer_index,
    is_pp_missing_parameter,
    make_empty_intermediate_tensors_factory,
@@ -49,6 +49,8 @@ from .utils import (
    maybe_prefix,
 )

+logger = init_logger(__name__)
+

 class ArcticMLP(nn.Module):
    def __init__(
@@ -382,7 +384,6 @@ class ArcticModel(nn.Module):
        cache_config = vllm_config.cache_config
        quant_config = vllm_config.quant_config

-        self.config = config
        self.vocab_size = config.vocab_size
        self.embed_tokens = VocabParallelEmbedding(
            self.vocab_size, config.hidden_size, org_num_embeddings=self.vocab_size
@@ -425,116 +426,6 @@ class ArcticModel(nn.Module):
        hidden_states = self.norm(hidden_states)
        return hidden_states

-    def load_weights(self, weights: Iterable[tuple[str, torch.Tensor]]) -> set[str]:
-        stacked_params_mapping = [
-            # (param_name, shard_name, shard_id)
-            ("qkv_proj", "q_proj", "q"),
-            ("qkv_proj", "k_proj", "k"),
-            ("qkv_proj", "v_proj", "v"),
-        ]
-
-        mlp_params_mapping: list[tuple[str, str, int]] = []
-        expert_params_mapping: list[tuple[str, str, int]] = []
-
-        for layer in range(self.config.num_hidden_layers):
-            is_moe_layer = (layer + 1) % self.config.moe_layer_frequency == 0
-            if is_moe_layer and self.config.use_residual:
-                mlp_params_mapping.append(
-                    (
-                        f"layers.{layer}.residual_mlp.w13.weight",
-                        f"layers.{layer}.residual_mlp.w1.weight",
-                        0,
-                    )
-                )
-                mlp_params_mapping.append(
-                    (
-                        f"layers.{layer}.residual_mlp.w13.weight",
-                        f"layers.{layer}.residual_mlp.w3.weight",
-                        1,
-                    )
-                )
-
-            if is_moe_layer:
-                for expert_id in range(self.config.num_local_experts):
-                    expert_params_mapping.append(
-                        ("ws", f"experts.{expert_id}.w1.weight", expert_id)
-                    )
-                    expert_params_mapping.append(
-                        ("w2s", f"experts.{expert_id}.w2.weight", expert_id)
-                    )
-                    expert_params_mapping.append(
-                        ("ws", f"experts.{expert_id}.w3.weight", expert_id)
-                    )
-            else:
-                mlp_params_mapping.append(
-                    (
-                        f"layers.{layer}.block_sparse_moe.mlp.w13.weight",
-                        f"layers.{layer}.block_sparse_moe.mlp.w1.weight",
-                        0,
-                    )
-                )
-                mlp_params_mapping.append(
-                    (
-                        f"layers.{layer}.block_sparse_moe.mlp.w13.weight",
-                        f"layers.{layer}.block_sparse_moe.mlp.w3.weight",
-                        1,
-                    )
-                )
-
-        params_dict = dict(self.named_parameters())
-        loaded_params: set[str] = set()
-
-        for name, loaded_weight in weights:
-            for param_name, weight_name, shard_id in stacked_params_mapping:
-                if weight_name not in name:
-                    continue
-                name = name.replace(weight_name, param_name)
-                # Skip loading extra bias for GPTQ models.
-                if name.endswith(".bias") and name not in params_dict:
-                    continue
-                if is_pp_missing_parameter(name, self):
-                    continue
-                param = params_dict[name]
-                weight_loader = param.weight_loader
-                weight_loader(param, loaded_weight, shard_id)
-                break
-            else:
-                for param_name, weight_name, shard_id in mlp_params_mapping:
-                    if weight_name not in name:
-                        continue
-                    name = name.replace(weight_name, param_name)
-                    if is_pp_missing_parameter(name, self):
-                        continue
-                    param = params_dict[name]
-                    weight_loader = param.weight_loader
-                    weight_loader(param, loaded_weight, shard_id)
-                    break
-                else:
-                    for param_name, weight_name, shard_id in expert_params_mapping:
-                        if weight_name not in name:
-                            continue
-                        name = name.replace(weight_name, param_name)
-                        if is_pp_missing_parameter(name, self):
-                            continue
-                        param = params_dict[name]
-                        weight_loader = param.weight_loader
-                        weight_loader(
-                            param, loaded_weight, weight_name, expert_id=shard_id
-                        )
-                        break
-                    else:
-                        if name.endswith(".bias") and name not in params_dict:
-                            continue
-                        if is_pp_missing_parameter(name, self):
-                            continue
-                        param = params_dict[name]
-                        weight_loader = getattr(
-                            param, "weight_loader", default_weight_loader
-                        )
-                        weight_loader(param, loaded_weight)
-            loaded_params.add(name)
-        return loaded_params
-

 class ArcticForCausalLM(nn.Module, SupportsPP, SupportsQuant):
    packed_modules_mapping = {"qkv_proj": ["q_proj", "k_proj", "v_proj"]}
@@ -587,8 +478,117 @@ class ArcticForCausalLM(nn.Module, SupportsPP, SupportsQuant):
        return logits

    def load_weights(self, weights: Iterable[tuple[str, torch.Tensor]]) -> set[str]:
-        loader = AutoWeightsLoader(
-            self,
-            skip_prefixes=(["lm_head."] if self.config.tie_word_embeddings else None),
+        stacked_params_mapping = [
+            # (param_name, shard_name, shard_id)
+            ("qkv_proj", "q_proj", "q"),
+            ("qkv_proj", "k_proj", "k"),
+            ("qkv_proj", "v_proj", "v"),
+        ]
+
+        mlp_params_mapping: list[tuple[str, str, int]] = []
+        expert_params_mapping: list[tuple[str, str, int]] = []
+        num_layers = self.config.num_hidden_layers
+
+        for layer in range(num_layers):
+            mlp_params_mapping.append(
+                (
+                    f"layers.{layer}.residual_mlp.w13.weight",
+                    f"layers.{layer}.residual_mlp.w1.weight",
+                    0,
+                )
+            )
+            mlp_params_mapping.append(
+                (
+                    f"layers.{layer}.residual_mlp.w13.weight",
+                    f"layers.{layer}.residual_mlp.w3.weight",
+                    1,
+                )
+            )
+            if layer % 2 == 0:
+                # MLP layers
+                mlp_params_mapping.append(
+                    (
+                        f"layers.{layer}.block_sparse_moe.mlp.w13.weight",
+                        f"layers.{layer}.block_sparse_moe.mlp.w1.weight",
+                        0,
+                    )
+                )
+                mlp_params_mapping.append(
+                    (
+                        f"layers.{layer}.block_sparse_moe.mlp.w13.weight",
+                        f"layers.{layer}.block_sparse_moe.mlp.w3.weight",
+                        1,
+                    )
+                )
+            else:
+                # MoE layers
+                for expert_id in range(self.config.num_local_experts):
+                    expert_params_mapping.append(
+                        ("ws", f"experts.{expert_id}.w1.weight", expert_id)
+                    )
+                    expert_params_mapping.append(
+                        ("w2s", f"experts.{expert_id}.w2.weight", expert_id)
+                    )
+                    expert_params_mapping.append(
+                        ("ws", f"experts.{expert_id}.w3.weight", expert_id)
+                    )
+
+        params_dict = dict(self.named_parameters())
+        loaded_params: set[str] = set()
+
+        logger.info(
+            "It will take ~10 minutes loading from the 16-bit weights. "
+            "Alternatively, use the prequantized 8-bit weights of arctic "
+            "and set load-format to `sharded_state` will accelerate loading."
        )
-        return loader.load_weights(weights)
+        for name, loaded_weight in weights:
+            for param_name, weight_name, shard_id in stacked_params_mapping:
+                if weight_name not in name:
+                    continue
+                name = name.replace(weight_name, param_name)
+                # Skip loading extra bias for GPTQ models.
+                if name.endswith(".bias") and name not in params_dict:
+                    continue
+                if is_pp_missing_parameter(name, self):
+                    continue
+                param = params_dict[name]
+                weight_loader = param.weight_loader
+                weight_loader(param, loaded_weight, shard_id)
+                break
+            else:
+                for param_name, weight_name, shard_id in mlp_params_mapping:
+                    if weight_name not in name:
+                        continue
+                    name = name.replace(weight_name, param_name)
+                    if is_pp_missing_parameter(name, self):
+                        continue
+                    param = params_dict[name]
+                    weight_loader = param.weight_loader
+                    weight_loader(param, loaded_weight, shard_id)
+                    break
+                else:
+                    for param_name, weight_name, shard_id in expert_params_mapping:
+                        if weight_name not in name:
+                            continue
+                        name = name.replace(weight_name, param_name)
+                        if is_pp_missing_parameter(name, self):
+                            continue
+                        param = params_dict[name]
+                        weight_loader = param.weight_loader
+                        weight_loader(
+                            param, loaded_weight, weight_name, expert_id=shard_id
+                        )
+                        break
+                    else:
+                        if name.endswith(".bias") and name not in params_dict:
+                            continue
+                        if is_pp_missing_parameter(name, self):
+                            continue
+                        param = params_dict[name]
+
+                        weight_loader = getattr(
+                            param, "weight_loader", default_weight_loader
+                        )
+                        weight_loader(param, loaded_weight)
+            loaded_params.add(name)
+        return loaded_params
@@ -184,16 +184,11 @@ class DeepSeekMTP(nn.Module, DeepseekV2MixtureOfExperts):
    def __init__(self, *, vllm_config: VllmConfig, prefix: str = ""):
        super().__init__()
        self.config = vllm_config.model_config.hf_config
-        self.quant_config = vllm_config.quant_config
        self.model = DeepSeekMultiTokenPredictor(
            vllm_config=vllm_config, prefix=maybe_prefix(prefix, "model")
        )
        # Set MoE hyperparameters
        self.set_moe_parameters()
-        self.is_fp4_ckpt = (
-            self.quant_config is not None
-            and self.quant_config.get_name() == "modelopt_fp4"
-        )

    def set_moe_parameters(self):
        self.expert_weights = []
@@ -246,16 +241,11 @@ class DeepSeekMTP(nn.Module, DeepseekV2MixtureOfExperts):
            ("gate_up_proj", "up_proj", 1),
            ("fused_qkv_a_proj", "q_a_proj", 0),
            ("fused_qkv_a_proj", "kv_a_proj_with_mqa", 1),
+            # Fused indexer wk + weights_proj
+            ("wk_weights_proj", "wk", 0),
+            ("wk_weights_proj", "weights_proj", 1),
        ]

-        if self.is_fp4_ckpt:
-            # Fused indexer wk + weights_proj (shard 0 = wk, shard 1 = weights_proj)
-            indexer_fused_mapping = [
-                ("wk_weights_proj", "wk", 0),
-                ("wk_weights_proj", "weights_proj", 1),
-            ]
-            stacked_params_mapping.extend(indexer_fused_mapping)
-
        expert_params_mapping = SharedFusedMoE.make_expert_params_mapping(
            self,
            ckpt_gate_proj_name="gate_proj",
@@ -625,11 +625,6 @@ class Indexer(nn.Module):
        super().__init__()
        self.vllm_config = vllm_config
        self.config = config
-        self.quant_config = quant_config
-        self.is_fp4_ckpt = (
-            self.quant_config is not None
-            and self.quant_config.get_name() == "modelopt_fp4"
-        )
        # self.indexer_cfg = config.attn_module_list_cfg[0]["attn_index"]
        self.topk_tokens = config.index_topk
        self.n_head = config.index_n_heads  # 64
@@ -644,36 +639,18 @@ class Indexer(nn.Module):
            quant_config=quant_config,
            prefix=f"{prefix}.wq_b",
        )
-        if self.is_fp4_ckpt:
-            # Fused wk + weights_proj: single GEMM producing [head_dim + n_head].
-            # weights_proj does not get quantized,
-            # so we run both with quant_config=None
-            # wk may be upcasted from the default quant;
-            # experiments show fusion is always faster unless WK proj is in FP4,
-            # which is not the case for all known quants.
-            self.wk_weights_proj = MergedColumnParallelLinear(
-                hidden_size,
-                [self.head_dim, self.n_head],
-                bias=False,
-                quant_config=None,
-                disable_tp=True,
-                prefix=f"{prefix}.wk_weights_proj",
-            )
-        else:
-            self.wk = ReplicatedLinear(
-                hidden_size,
-                self.head_dim,
-                bias=False,
-                quant_config=quant_config,
-                prefix=f"{prefix}.wk",
-            )
-            self.weights_proj = ReplicatedLinear(
-                hidden_size,
-                self.n_head,
-                bias=False,
-                quant_config=None,
-                prefix=f"{prefix}.weights_proj",
-            )
+        # Fused wk + weights_proj: single GEMM producing [head_dim + n_head].
+        # weights_proj does not get quantized, so we run both with quant_config=None
+        # wk may be upcasted from the default quant; experiments show fusion is always
+        # faster unless WK proj is in FP4, which is not the case for all known quants.
+        self.wk_weights_proj = MergedColumnParallelLinear(
+            hidden_size,
+            [self.head_dim, self.n_head],
+            bias=False,
+            quant_config=None,
+            disable_tp=True,
+            prefix=f"{prefix}.wk_weights_proj",
+        )
        self.k_norm = LayerNorm(self.head_dim, eps=1e-6)
        self.softmax_scale = self.head_dim**-0.5

@@ -714,14 +691,11 @@ class Indexer(nn.Module):
        q_pe, q_nope = torch.split(
            q, [self.rope_dim, self.head_dim - self.rope_dim], dim=-1
        )
-        if self.is_fp4_ckpt:
-            # Fused wk + weights_proj: one GEMM, then split
-            kw, _ = self.wk_weights_proj(hidden_states)
-            k = kw[:, : self.head_dim]
-            weights = kw[:, self.head_dim :]
-        else:
-            k, _ = self.wk(hidden_states)
-            weights, _ = self.weights_proj(hidden_states)
+
+        # Fused wk + weights_proj: one GEMM, then split
+        kw, _ = self.wk_weights_proj(hidden_states)
+        k = kw[:, : self.head_dim]
+        weights_raw = kw[:, self.head_dim :]

        k = self.k_norm(k)
        k_pe, k_nope = torch.split(
@@ -752,7 +726,7 @@ class Indexer(nn.Module):
        q_scale = q_scale.view(-1, self.n_head, 1)

        weights = (
-            weights.unsqueeze(-1) * q_scale * self.softmax_scale * self.n_head**-0.5
+            weights_raw.unsqueeze(-1) * q_scale * self.softmax_scale * self.n_head**-0.5
        )
        weights = weights.squeeze(-1)

@@ -1340,10 +1314,6 @@ class DeepseekV2ForCausalLM(
        quant_config = vllm_config.quant_config
        self.config = config
        self.quant_config = quant_config
-        self.is_fp4_ckpt = (
-            self.quant_config is not None
-            and self.quant_config.get_name() == "modelopt_fp4"
-        )

        qk_nope_head_dim = getattr(config, "qk_nope_head_dim", 0)
        qk_rope_head_dim = getattr(config, "qk_rope_head_dim", 0)
@@ -1469,13 +1439,12 @@ class DeepseekV2ForCausalLM(
            ("qkv_proj", "k_proj", "k"),
            ("qkv_proj", "v_proj", "v"),
        ]
-        if self.is_fp4_ckpt:
-            # Fused indexer wk + weights_proj (shard 0 = wk, shard 1 = weights_proj)
-            indexer_fused_mapping = [
-                ("wk_weights_proj", "wk", 0),
-                ("wk_weights_proj", "weights_proj", 1),
-            ]
-            stacked_params_mapping.extend(indexer_fused_mapping)
+        # Fused indexer wk + weights_proj (shard 0 = wk, shard 1 = weights_proj)
+        indexer_fused_mapping = [
+            ("wk_weights_proj", "wk", 0),
+            ("wk_weights_proj", "weights_proj", 1),
+        ]
+        stacked_params_mapping.extend(indexer_fused_mapping)

        if self.use_mha:
            stacked_params_mapping.extend(mha_params_mapping)
@@ -620,7 +620,6 @@ class Qwen3_5ForConditionalGeneration(Qwen3VLForConditionalGeneration, IsHybrid)
            self.packed_modules_mapping = {k: list(v) for k, v in base.items()}
            self.packed_modules_mapping.pop("in_proj_qkvz", None)
            self.packed_modules_mapping["in_proj_qkv"] = ["in_proj_qkv"]
-            self.packed_modules_mapping["in_proj_z"] = ["in_proj_z"]

    def embed_input_ids(
        self,
@@ -381,12 +381,7 @@ class CpuPlatform(Platform):
    def get_global_cpu_mask(cls) -> set[int]:
        # get global cpu mask
        if cls.global_cpu_mask is None:
-            if hasattr(os, "sched_getaffinity"):
-                cls.global_cpu_mask = os.sched_getaffinity(0)
-            else:
-                # macOS does not support sched_getaffinity
-                cpu_count = os.cpu_count() or 1
-                cls.global_cpu_mask = set(range(cpu_count))
+            cls.global_cpu_mask = os.sched_getaffinity(0)
        return cls.global_cpu_mask

    @classmethod
@@ -207,15 +207,6 @@ class Platform:
        """
        return cls.simple_compile_backend

-    @classmethod
-    def import_ir_kernels(cls) -> None:
-        """
-        The default implementation imports ``vllm.kernels``, which registers
-        the built-in IR op implementations. Out-of-tree (OOT) platforms should
-        override this method to import their own kernel modules.
-        """
-        import vllm.kernels  # noqa: F401
-
    @classmethod
    def device_id_to_physical_device_id(cls, device_id: int):
        # Treat empty device control env var as unset. This is a valid
@@ -402,6 +402,7 @@ class RocmPlatform(Platform):
        "gguf",
        "quark",
        "mxfp4",
+        "petit_nvfp4",
        "torchao",
        "bitsandbytes",
    ]
@@ -7,7 +7,6 @@ Copyright (c) 2026 Cambridge Greys Ltd

 import json
 import os
-import platform
 import subprocess


@@ -39,37 +38,10 @@ def parse_mask(mask):
    return set(result)


-def _get_default_affinity() -> set[int]:
-    """Get the set of CPUs the process is allowed to run on."""
-    if hasattr(os, "sched_getaffinity"):
-        return os.sched_getaffinity(0)
-    # macOS does not support sched_getaffinity; fall back to cpu_count
-    cpu_count = os.cpu_count() or 1
-    return set(range(cpu_count))
-
-
-def _get_cpu_topology_json() -> bytes:
-    """Get CPU topology as JSON.
-
-    On Linux this uses ``lscpu -Je``.  On other platforms (e.g. macOS) we
-    synthesize a simple topology where every logical CPU is its own core
-    on NUMA node 0, which is sufficient for the OMP place-list builder.
-    """
-    if platform.system() == "Linux":
-        return subprocess.run(["lscpu", "-Je"], check=True, capture_output=True).stdout
-
-    # Fallback for non-Linux (macOS, etc.)
-    cpu_count = os.cpu_count() or 1
-    cpus = []
-    for i in range(cpu_count):
-        cpus.append({"cpu": str(i), "core": str(i), "node": "0"})
-    return json.dumps({"cpus": cpus}).encode()
-
-
 def enumerate_resources(resource_map, mask=None, allowed=None):
    """Enumerate system resources"""
    if allowed is None:
-        allowed = _get_default_affinity()
+        allowed = os.sched_getaffinity(0)
    if mask is not None:
        allowed = allowed & mask

@@ -168,7 +140,9 @@ class OMPProcessManager:
            else:
                masks = [None]
            if mock is None:
-                data = _get_cpu_topology_json()
+                data = subprocess.run(
+                    ["lscpu", "-Je"], check=True, capture_output=True
+                ).stdout
            else:
                with open(mock, mode="rb") as jf:
                    data = jf.read()
@@ -4192,7 +4192,6 @@ class GPUModelRunner(
        spec_config = self.speculative_config
        propose_drafts_after_bookkeeping = False
        if spec_config is not None:
-            # Decide whether to run the drafter or zero out draft tokens.
            input_fits_in_drafter = spec_decode_common_attn_metadata is not None and (
                spec_decode_common_attn_metadata.max_seq_len + self.num_spec_tokens
                <= self.effective_drafter_max_model_len
@@ -4228,6 +4227,10 @@ class GPUModelRunner(
                    self._copy_valid_sampled_token_count(
                        next_token_ids, valid_sampled_tokens_count
                    )
+                    self._draft_token_ids = torch.zeros(
+                        1, device=self.device, dtype=torch.int32
+                    ).expand(len(self.input_batch.req_ids), self.num_spec_tokens)
+                    self._copy_draft_token_ids_to_cpu(scheduler_output, zeros_only=True)
            elif (
                spec_config.use_ngram_gpu()
                and not spec_config.disable_padded_drafter_batch
@@ -4250,20 +4253,15 @@ class GPUModelRunner(
                    self._copy_valid_sampled_token_count(
                        next_token_ids, valid_sampled_tokens_count
                    )
+                    # Since we couldn't run the drafter,
+                    # just use zeros for the draft tokens.
+                    self._draft_token_ids = torch.zeros(
+                        1, device=self.device, dtype=torch.int32
+                    ).expand(len(self.input_batch.req_ids), self.num_spec_tokens)
+                    self._copy_draft_token_ids_to_cpu(scheduler_output, zeros_only=True)
            else:
                propose_drafts_after_bookkeeping = input_fits_in_drafter

-            if not input_fits_in_drafter:
-                # Zero out draft tokens so the scheduler doesn't schedule
-                # stale drafts from the previous step.
-                # For Nemotron-H: it is necessary to zero out the draft tokens,
-                # otherwise the stale tokens will corrupt Mamba recurrent
-                # state and logprobs for sequences near max_model_len.
-                self._draft_token_ids = torch.zeros(
-                    1, device=self.device, dtype=torch.int32
-                ).expand(len(self.input_batch.req_ids), self.num_spec_tokens)
-                self._copy_draft_token_ids_to_cpu(scheduler_output, zeros_only=True)
-
        with record_function_or_nullcontext("gpu_model_runner: bookkeep"):
            (
                num_nans_in_logits,
@@ -6846,7 +6844,7 @@ class GPUModelRunner(
            # group
            self.drafter.validate_same_kv_cache_group(kv_cache_config)

-        if has_kv_transfer_group() and not is_profiling:
+        if has_kv_transfer_group():
            kv_transfer_group = get_kv_transfer_group()
            if self.cross_layers_kv_cache is not None:
                assert self.cross_layers_attn_backend is not None
@@ -31,7 +31,7 @@ _manager: "WorkspaceManager | None" = None
 class WorkspaceManager:
    """Manager for workspace allocation.

-    Manages one workspace buffer per active ubatch slot.
+    Manages workspace buffers for DBO (Dual Batch Overlap) execution.
    Can be locked to prevent further growth during execution.
    """

@@ -39,9 +39,7 @@ class WorkspaceManager:
        self._device = device
        # Cache num ubatches at init based on configuration (default to 1)
        self._num_ubatches = num_ubatches if num_ubatches is not None else 1
-        self._current_workspaces: list[torch.Tensor | None] = [
-            None
-        ] * self._num_ubatches
+        self._current_workspaces: list[torch.Tensor | None] = [None, None]
        self._locked: bool = False

    @staticmethod
@@ -226,7 +224,7 @@ def init_workspace_manager(

    Args:
        device: The device to allocate workspace on.
-        num_ubatches: Number of workspace ubatch slots. Defaults to 1.
+        num_ubatches: Number of micro-batches. Defaults to 1.
    """
    global _manager
    if _manager is not None:
Author	SHA1	Message	Date
Wentao YeandGitHub	beefff2174	Merge branch 'main' into wentao-skip-work-when-empty	2026-04-03 11:31:23 -04:00
yewentao256	163266d0b2	refactor only inside nixl Signed-off-by: yewentao256 <zhyanwentao@126.com>	2026-04-01 18:28:13 +00:00
yewentao256	a2fd28a7e1	Merge branch 'main' into wentao-skip-work-when-empty Signed-off-by: yewentao256 <zhyanwentao@126.com>	2026-04-01 18:09:06 +00:00
Wentao YeandGitHub	7a80ac928f	Merge branch 'main' into wentao-skip-work-when-empty	2026-03-31 13:01:59 -04:00
yewentao256	0ed11013b4	check empty using kv connector's method Signed-off-by: yewentao256 <zhyanwentao@126.com>	2026-03-30 21:06:12 +00:00
yewentao256	6d568b995a	Merge branch 'main' into wentao-skip-work-when-empty	2026-03-30 19:17:12 +00:00
Wentao YeandGitHub	dfe9decbcb	Merge branch 'main' into wentao-skip-work-when-empty	2026-03-28 10:28:43 -04:00
Wentao YeandGitHub	1324e6ff67	Merge branch 'main' into wentao-skip-work-when-empty	2026-03-27 15:52:55 -04:00
Wentao YeandGitHub	c1aba6d7ae	Merge branch 'main' into wentao-skip-work-when-empty	2026-03-27 09:42:19 -04:00
Wentao Ye GitHubgemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com>	200bef28c9	Update vllm/v1/worker/gpu/kv_connector.py Co-authored-by: gemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com> Signed-off-by: Wentao Ye <44945378+yewentao256@users.noreply.github.com>	2026-03-26 18:04:12 -04:00
yewentao256	fd9820bbf9	skip kv connector empty work Signed-off-by: yewentao256 <zhyanwentao@126.com>	2026-03-26 21:50:28 +00:00