remove multiple dead codes

Signed-off-by: yewentao256 <zhyanwentao@126.com>

.buildkite/test-amd.yaml

@@ -1576,14 +1576,10 @@ steps:
   - vllm/third_party/flash_linear_attention/ops/kda.py
   - vllm/third_party/flash_linear_attention/ops/chunk_delta_h.py
   - vllm/third_party/flash_linear_attention/ops/l2norm.py
-  - vllm/models/kimi_k3/nvidia/kda.py
-  - vllm/models/kimi_k3/nvidia/kda_metadata.py
-  - vllm/models/kimi_k3/nvidia/ops/third_party/kda/
-  - tests/models/kimi_k3/test_kda.py
-  - tests/models/kimi_k3/test_kda_metadata.py
+  - tests/kernels/test_kda.py
   - vllm/platforms/rocm.py
   commands:
-  - pytest -v -s models/kimi_k3/test_kda.py models/kimi_k3/test_kda_metadata.py
+  - pytest -v -s kernels/test_kda.py
 
 - label: Kernels Mamba Test # TBD
   timeout_in_minutes: 180

.gitea/workflows/build-gb10.yml

@@ -1,105 +0,0 @@
-name: Build GB10 ARM64 image
-
-on:
-  workflow_dispatch:
-  push:
-    branches:
-      - karylab/gb10
-    paths:
-      - ".gitea/workflows/build-gb10.yml"
-      - "CMakeLists.txt"
-      - "cmake/**"
-      - "csrc/**"
-      - "docker/**"
-      - "pyproject.toml"
-      - "requirements/**"
-      - "setup.py"
-      - "vllm/**"
-
-env:
-  REGISTRY: git.karylab.com
-  IMAGE_NAME: karylab_agents/vllm-gb10
-
-jobs:
-  build-arm64:
-    # This label is registered by gitea-act-runner.yml.
-    runs-on: vllm-x86-builder
-    timeout-minutes: 720
-
-    # Do not add pull_request as a trigger. This runner controls the build
-    # host's Docker socket and must only execute reviewed code on main.
-    permissions:
-      contents: read
-
-    steps:
-      - name: Check out source
-        uses: actions/checkout@v4
-
-      - name: Enable ARM64 emulation
-        uses: docker/setup-qemu-action@v3
-        with:
-          platforms: arm64
-
-      - name: Create Buildx builder
-        uses: docker/setup-buildx-action@v3
-
-      - name: Validate registry credential
-        env:
-          REGISTRY_TOKEN: ${{ secrets.REGISTRY_TOKEN }}
-        run: |
-          if [ -z "${REGISTRY_TOKEN}" ]; then
-            echo "Missing repository Actions secret: REGISTRY_TOKEN" >&2
-            exit 1
-          fi
-
-      - name: Log in to Gitea Container Registry
-        env:
-          REGISTRY_TOKEN: ${{ secrets.REGISTRY_TOKEN }}
-        run: |
-          printf '%s' "${REGISTRY_TOKEN}" |
-            docker login "${REGISTRY}" \
-              --username karylab_agents \
-              --password-stdin
-
-      - name: Derive immutable image tag
-        id: image
-        shell: bash
-        run: |
-          short_sha="$(git rev-parse --short=12 HEAD)"
-          echo "tag=gb10-${short_sha}" >> "${GITHUB_OUTPUT}"
-          echo "revision=$(git rev-parse HEAD)" >> "${GITHUB_OUTPUT}"
-
-      - name: Build and push ARM64 image
-        env:
-          IMAGE_TAG: ${{ steps.image.outputs.tag }}
-          REVISION: ${{ steps.image.outputs.revision }}
-        run: |
-          image="${REGISTRY}/${IMAGE_NAME}"
-          run_url="${{ github.server_url }}/${{ github.repository }}/actions/runs/${{ github.run_id }}"
-
-          docker buildx build \
-            --platform linux/arm64 \
-            --file docker/Dockerfile \
-            --target vllm-openai \
-            --build-arg "torch_cuda_arch_list=12.0" \
-            --build-arg "max_jobs=4" \
-            --build-arg "nvcc_threads=8" \
-            --build-arg "VLLM_BUILD_COMMIT=${REVISION}" \
-            --build-arg "VLLM_BUILD_PIPELINE=gitea-actions" \
-            --build-arg "VLLM_BUILD_URL=${run_url}" \
-            --build-arg "VLLM_IMAGE_TAG=${image}:${IMAGE_TAG}" \
-            --cache-from "type=registry,ref=${image}:buildcache" \
-            --cache-to "type=registry,ref=${image}:buildcache,mode=max,compression=zstd,oci-mediatypes=true" \
-            --label "org.opencontainers.image.revision=${REVISION}" \
-            --label "org.opencontainers.image.source=${{ github.server_url }}/${{ github.repository }}" \
-            --tag "${image}:${IMAGE_TAG}" \
-            --tag "${image}:latest" \
-            --push \
-            .
-
-      - name: Verify published image
-        env:
-          IMAGE_TAG: ${{ steps.image.outputs.tag }}
-        run: |
-          docker buildx imagetools inspect \
-            "${REGISTRY}/${IMAGE_NAME}:${IMAGE_TAG}"

tests/model_executor/test_deepseek_v4_mhc_warmup.py

@@ -1,668 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-"""CPU-pure tests for DeepSeek V4 MHC warmup selection and gating logic.
-
-The TileLang JIT kernels require CUDA; these tests verify CPU-side selection,
-layer-finding, and gating without a GPU.  ``_select_mhc_split_key_token_sizes``
-tests monkeypatch ``compute_num_split`` (CUDA-backed) with a deterministic fake
-so the real function runs under CPU-pure coverage.
-"""
-
-import importlib
-import sys
-import types
-from types import SimpleNamespace
-
-import pytest
-import torch
-
-from vllm.model_executor.warmup.deepseek_v4_mhc_warmup import (
-    _find_deepseek_v4_model,
-    _find_first_mhc_layer,
-    _normalize_token_sizes,
-    _select_mhc_split_key_token_sizes,
-    _select_mhc_warmup_token_sizes,
-)
-from vllm.utils.math_utils import cdiv
-
-# ── Helpers ───────────────────────────────────────────────────────────────
-
-
-def _fake_compute_num_split(n_sms: int):
-    """Deterministic ``compute_num_split`` with explicit SM count instead of
-    ``torch.cuda``."""
-
-    def _inner(block_k: int, k: int | None, grid_size: int) -> int:
-        split_k = n_sms // grid_size
-        if k is not None:
-            split_k = min(split_k, cdiv(k, block_k) // 4)
-        return max(split_k, 1)
-
-    return _inner
-
-
-def _patch_compute_split(*, n_sms: int, monkeypatch: pytest.MonkeyPatch) -> None:
-    """Stub ``tilelang_kernels`` in ``sys.modules`` via ``monkeypatch.setitem``
-    so the local import inside ``_select_mhc_split_key_token_sizes`` resolves
-    without triggering the real module (which requires TileLang + CUDA).  The
-    stub is automatically restored after each test."""
-
-    MODULE_PATH = "vllm.model_executor.kernels.mhc.tilelang_kernels"
-
-    # Ensure parent packages exist in sys.modules so the dotted-path import
-    # resolves through parent lookups.
-    for parent_path in (
-        "vllm",
-        "vllm.model_executor",
-        "vllm.model_executor.kernels",
-        "vllm.model_executor.kernels.mhc",
-    ):
-        if parent_path not in sys.modules:
-            monkeypatch.setitem(sys.modules, parent_path, types.ModuleType(parent_path))
-
-    stub = types.ModuleType(MODULE_PATH)
-    stub.compute_num_split = _fake_compute_num_split(n_sms)
-    monkeypatch.setitem(sys.modules, MODULE_PATH, stub)
-
-
-# ── _normalize_token_sizes ───────────────────────────────────────────────
-
-
-class TestNormalizeTokenSizes:
-    def test_empty_when_no_sizes(self) -> None:
-        assert _normalize_token_sizes((), max_tokens=100) == []
-
-    def test_removes_out_of_range(self) -> None:
-        assert _normalize_token_sizes([0, 1, 50, 100, 200], max_tokens=100) == [
-            1,
-            50,
-            100,
-        ]
-
-    def test_deduplicates_and_sorts(self) -> None:
-        assert _normalize_token_sizes([4, 1, 4, 8, 2], max_tokens=100) == [1, 2, 4, 8]
-
-    def test_accepts_iterator(self) -> None:
-        assert _normalize_token_sizes(iter({1, 2, 3}), max_tokens=10) == [1, 2, 3]
-
-
-# ── _select_mhc_warmup_token_sizes ────────────────────────────────────────
-
-
-class TestSelectMhcWarmupTokenSizes:
-    def test_empty_on_zero_max_tokens(self) -> None:
-        assert (
-            _select_mhc_warmup_token_sizes(max_tokens=0, cudagraph_capture_sizes=[])
-            == []
-        )
-
-    def test_contains_1(self) -> None:
-        assert 1 in _select_mhc_warmup_token_sizes(
-            max_tokens=10, cudagraph_capture_sizes=[]
-        )
-
-    def test_bounded_by_max_tokens(self) -> None:
-        sizes = _select_mhc_warmup_token_sizes(max_tokens=5, cudagraph_capture_sizes=[])
-        assert all(1 <= s <= 5 for s in sizes)
-
-    def test_includes_cudagraph_capture_sizes(self) -> None:
-        sizes = _select_mhc_warmup_token_sizes(
-            max_tokens=100, cudagraph_capture_sizes=[7, 33]
-        )
-        assert 7 in sizes and 33 in sizes
-
-    def test_includes_max_auto_tokens(self) -> None:
-        sizes = _select_mhc_warmup_token_sizes(
-            max_tokens=100, cudagraph_capture_sizes=[]
-        )
-        assert 100 in sizes
-
-    def test_respects_auto_warmup_cap(self) -> None:
-        sizes = _select_mhc_warmup_token_sizes(
-            max_tokens=20000, cudagraph_capture_sizes=[]
-        )
-        assert max(sizes) == 16384
-
-    def test_does_not_include_zero(self) -> None:
-        assert _select_mhc_warmup_token_sizes(
-            max_tokens=1, cudagraph_capture_sizes=[]
-        ) == [1]
-
-
-# ── _select_mhc_split_key_token_sizes ─────────────────────────────────────
-
-
-class TestSelectMhcSplitKeyTokenSizes:
-    """Real ``_select_mhc_split_key_token_sizes`` invoked under a deterministic
-    ``compute_num_split`` monkeypatch.  No CUDA required outside the skip-guard
-    cross-check below."""
-
-    # ── Broadcast-variant key counts (k_size = hidden_size) ──────────────
-
-    @pytest.mark.parametrize(
-        "max_tokens,k_size,n_sms,expected_keys",
-        [
-            (8192, 4096, 188, 16),  # RTX PRO 6000 Blackwell, broadcast K
-            (8192, 4096, 132, 15),  # H100 SXM
-            (8192, 4096, 80, 12),
-            (8192, 2048, 188, 8),
-            (8192, 1024, 188, 4),
-            (1024, 4096, 188, 6),
-            (1, 4096, 188, 1),
-        ],
-    )
-    def test_broadcast_key_count(
-        self, monkeypatch, max_tokens, k_size, n_sms, expected_keys
-    ) -> None:
-        _patch_compute_split(n_sms=n_sms, monkeypatch=monkeypatch)
-        reps = _select_mhc_split_key_token_sizes(max_tokens=max_tokens, k_size=k_size)
-        assert len(reps) == expected_keys, (
-            f"Expected {expected_keys} keys for max_tokens={max_tokens}, "
-            f"k_size={k_size}, n_sms={n_sms}, got {len(reps)}"
-        )
-
-    # ── Non-broadcast-variant key counts (k_size = hc_mult * hidden_size) ──
-
-    @pytest.mark.parametrize(
-        "max_tokens,k_size,n_sms,expected_keys",
-        [
-            (8192, 65536, 188, 26),  # DSv4 default: hc_mult=16, hidden=4096
-            (8192, 65536, 132, 22),  # H100 SXM
-            (8192, 65536, 80, 16),
-            (8192, 32768, 188, 26),  # hc_mult=8, hidden=4096; same as 188 SM full
-            (8192, 8192, 188, 22),
-            (1024, 65536, 188, 16),
-            (1, 65536, 188, 1),
-        ],
-    )
-    def test_non_broadcast_key_count(
-        self, monkeypatch, max_tokens, k_size, n_sms, expected_keys
-    ) -> None:
-        _patch_compute_split(n_sms=n_sms, monkeypatch=monkeypatch)
-        reps = _select_mhc_split_key_token_sizes(max_tokens=max_tokens, k_size=k_size)
-        assert len(reps) == expected_keys, (
-            f"Expected {expected_keys} keys for max_tokens={max_tokens}, "
-            f"k_size={k_size}, n_sms={n_sms}, got {len(reps)}"
-        )
-
-    # ── Semantic invariants ─────────────────────────────────────────────
-
-    def test_all_keys_distinct(self, monkeypatch) -> None:
-        _patch_compute_split(n_sms=188, monkeypatch=monkeypatch)
-        reps = _select_mhc_split_key_token_sizes(max_tokens=8192, k_size=4096)
-        fake_cns = _fake_compute_num_split(188)
-        ns_values = [fake_cns(64, 4096, cdiv(t, 64)) for t in reps]
-        assert len(set(ns_values)) == len(ns_values), (
-            f"Duplicate n_splits: {dict(zip(reps, ns_values))}"
-        )
-
-    def test_last_key_n_splits_is_one(self, monkeypatch) -> None:
-        _patch_compute_split(n_sms=188, monkeypatch=monkeypatch)
-        reps = _select_mhc_split_key_token_sizes(max_tokens=8192, k_size=4096)
-        last_ns = _fake_compute_num_split(188)(64, 4096, cdiv(reps[-1], 64))
-        assert last_ns == 1, f"Last key n_splits={last_ns}, expected 1"
-
-    # ── Exact broadcast sequence (k_size=hidden_size=4096, 188 SMs) ──────
-
-    def test_exact_broadcast_sequence_188_sms(self, monkeypatch) -> None:
-        # fmt: off
-        expected = [1, 705, 769, 833, 897, 961, 1089, 1153,
-                    1281, 1473, 1665, 1985, 2369, 3009, 3969, 6017]
-        # fmt: on
-        _patch_compute_split(n_sms=188, monkeypatch=monkeypatch)
-        reps = _select_mhc_split_key_token_sizes(max_tokens=8192, k_size=4096)
-        assert reps == expected, (
-            f"188-SM broadcast (K=4096) sequence mismatch\n"
-            f"  Expected ({len(expected)}): {expected}\n"
-            f"  Got      ({len(reps)}): {reps}"
-        )
-
-    # ── Exact non-broadcast sequence (k_size=65536, 188 SMs) ─────────────
-
-    def test_exact_non_broadcast_sequence_188_sms(self, monkeypatch) -> None:
-        # fmt: off
-        expected = [1, 65, 129, 193, 257, 321, 385, 449, 513,
-                    577, 641, 705, 769, 833, 897, 961, 1089,
-                    1153, 1281, 1473, 1665, 1985, 2369, 3009,
-                    3969, 6017]
-        # fmt: on
-        _patch_compute_split(n_sms=188, monkeypatch=monkeypatch)
-        reps = _select_mhc_split_key_token_sizes(max_tokens=8192, k_size=65536)
-        assert reps == expected, (
-            f"188-SM non-broadcast (K=65536) sequence mismatch\n"
-            f"  Expected ({len(expected)}): {expected}\n"
-            f"  Got      ({len(reps)}): {reps}"
-        )
-
-    # ── CUDA real cross-check (skipped when GPU or TileLang unavailable) ─
-
-    def test_selector_matches_compute_num_split(self, monkeypatch) -> None:
-        if not torch.cuda.is_available():
-            pytest.skip("CUDA required for production cross-check")
-
-        MODULE_PATH = "vllm.model_executor.kernels.mhc.tilelang_kernels"
-        monkeypatch.delitem(sys.modules, MODULE_PATH, raising=False)
-        importlib.invalidate_caches()
-
-        try:
-            from vllm.model_executor.kernels.mhc.tilelang_kernels import (
-                compute_num_split,
-            )
-        except Exception as e:
-            pytest.skip(f"TileLang real module unavailable: {e}")
-
-        real_module = sys.modules[MODULE_PATH]
-        assert real_module.__file__ is not None
-        assert "vllm" in real_module.__file__
-
-        reps = _select_mhc_split_key_token_sizes(max_tokens=8192, k_size=4096)
-        assert len(reps) >= 1
-        ns_values = [compute_num_split(64, 4096, cdiv(t, 64)) for t in reps]
-        assert len(set(ns_values)) == len(ns_values), (
-            f"Duplicate n_splits: {dict(zip(reps, ns_values))}"
-        )
-        assert all(ns >= 1 for ns in ns_values), "Some n_splits are zero"
-
-
-# ── _find_first_mhc_layer ───────────────────────────────────────────────
-
-
-class TestFindFirstMhcLayer:
-    def test_finds_layer_with_all_required_attrs(self) -> None:
-        class MockLayer:
-            hc_pre = hc_post = hc_attn_fn = hc_attn_scale = None
-            hc_attn_base = hc_ffn_fn = hc_ffn_scale = hc_ffn_base = None
-
-        MockLayer.__name__ = MockLayer.__qualname__ = "DeepseekV4DecoderLayer"
-
-        class MockModel:
-            def modules(self):
-                yield self
-                yield MockLayer()
-
-        result = _find_first_mhc_layer(MockModel())
-        assert result is not None
-        assert result.__class__.__name__ == "DeepseekV4DecoderLayer"
-
-    def test_skips_layer_missing_required_attr(self) -> None:
-        class IncompleteLayer:
-            hc_post = hc_attn_fn = hc_attn_scale = hc_attn_base = None
-            hc_ffn_fn = hc_ffn_scale = hc_ffn_base = None
-
-        IncompleteLayer.__name__ = IncompleteLayer.__qualname__ = (
-            "DeepseekV4DecoderLayer"
-        )
-
-        class MockModel:
-            def modules(self):
-                yield self
-                yield IncompleteLayer()
-
-        assert _find_first_mhc_layer(MockModel()) is None
-
-
-# ── _find_deepseek_v4_model ──────────────────────────────────────────────
-
-
-class TestFindDeepseekV4Model:
-    def test_finds_model_with_all_required_attrs(self) -> None:
-        class MockDsModel:
-            hc_head_fn = hc_head_scale = hc_head_base = None
-
-        MockDsModel.__name__ = MockDsModel.__qualname__ = "DeepseekV4Model"
-
-        class MockModel:
-            def modules(self):
-                yield self
-                yield MockDsModel()
-
-        result = _find_deepseek_v4_model(MockModel())
-        assert result is not None
-        assert result.__class__.__name__ == "DeepseekV4Model"
-
-    def test_skips_model_missing_required_attr(self) -> None:
-        class IncompleteModel:
-            hc_head_fn = hc_head_scale = None
-
-        IncompleteModel.__name__ = IncompleteModel.__qualname__ = "DeepseekV4Model"
-
-        class MockModel:
-            def modules(self):
-                yield self
-                yield IncompleteModel()
-
-        assert _find_deepseek_v4_model(MockModel()) is None
-
-    def test_skips_wrong_class_name(self) -> None:
-        class OtherModel:
-            hc_head_fn = hc_head_scale = hc_head_base = None
-
-        OtherModel.__name__ = OtherModel.__qualname__ = "OtherModel"
-
-        class MockModel:
-            def modules(self):
-                yield self
-                yield OtherModel()
-
-        assert _find_deepseek_v4_model(MockModel()) is None
-
-
-# ── Broadcast no-op gates ─────────────────────────────────────────────────
-
-
-class TestWarmupBroadcastNoOpConditions:
-    @staticmethod
-    def _install_broadcast_spy(monkeypatch) -> list[str]:
-        calls: list[str] = []
-        module_path = "vllm.model_executor.kernels.mhc.tilelang"
-        stub = types.ModuleType(module_path)
-        stub.mhc_pre_broadcast_tilelang = lambda *args, **kwargs: calls.append(
-            "mhc_pre_broadcast_tilelang"
-        )
-        monkeypatch.setitem(sys.modules, module_path, stub)
-        return calls
-
-    @staticmethod
-    def _model_with_layer(layer):
-        layer.__name__ = layer.__qualname__ = "DeepseekV4DecoderLayer"
-
-        class MockModel:
-            def modules(self):
-                yield self
-                yield layer()
-
-        return MockModel()
-
-    def test_noop_when_no_broadcast_layer(self) -> None:
-        from vllm.model_executor.warmup.deepseek_v4_mhc_warmup import (
-            _warmup_broadcast_mhc,
-        )
-
-        class NoBroadcastLayer:
-            hc_attn_fn_broadcast = None
-
-        _warmup_broadcast_mhc(
-            self._model_with_layer(NoBroadcastLayer), token_sizes=[1, 2, 4]
-        )
-
-    def test_noop_when_device_not_cuda(self, monkeypatch) -> None:
-        from vllm.model_executor.warmup.deepseek_v4_mhc_warmup import (
-            _warmup_broadcast_mhc,
-        )
-
-        calls = self._install_broadcast_spy(monkeypatch)
-
-        class CpuLayer:
-            hc_attn_fn_broadcast = torch.empty(0)
-            hc_attn_fn = torch.empty(0, device="cpu")
-
-        _warmup_broadcast_mhc(self._model_with_layer(CpuLayer), token_sizes=[1, 2, 4])
-        assert not calls
-
-    def test_noop_when_broadcast_not_a_tensor(self, monkeypatch) -> None:
-        from vllm.model_executor.warmup.deepseek_v4_mhc_warmup import (
-            _warmup_broadcast_mhc,
-        )
-
-        calls = self._install_broadcast_spy(monkeypatch)
-
-        class BoolBroadcastLayer:
-            hc_attn_fn_broadcast = True
-            hc_attn_fn = SimpleNamespace(device=torch.device("cuda"))
-
-        _warmup_broadcast_mhc(
-            self._model_with_layer(BoolBroadcastLayer), token_sizes=[1, 2, 4]
-        )
-        assert not calls
-
-    def test_noop_when_broadcast_device_mismatch(self, monkeypatch) -> None:
-        from unittest import mock
-
-        from vllm.model_executor.warmup.deepseek_v4_mhc_warmup import (
-            _warmup_broadcast_mhc,
-        )
-
-        calls = self._install_broadcast_spy(monkeypatch)
-        broadcast = mock.MagicMock(spec=torch.Tensor)
-        broadcast.device = torch.device("cuda:1")
-
-        class MismatchedLayer:
-            hc_attn_fn_broadcast = broadcast
-            hc_attn_fn = SimpleNamespace(device=torch.device("cuda:0"))
-
-        _warmup_broadcast_mhc(
-            self._model_with_layer(MismatchedLayer), token_sizes=[1, 2, 4]
-        )
-        assert not calls
-
-
-# ── _warmup_layer_mhc union behavior ─────────────────────────────────────
-
-
-class TestWarmupLayerMhcUnion:
-    """Verify ``_warmup_layer_mhc`` unions general ``token_sizes`` with
-    split-key reps computed from ``hc_mult * hidden_size``, and calls
-    ``hc_pre``/``hc_post`` for each token size in the union."""
-
-    def test_unions_general_sizes_with_split_key_reps(self, monkeypatch) -> None:
-        from vllm.model_executor.warmup.deepseek_v4_mhc_warmup import (
-            _warmup_layer_mhc,
-        )
-
-        # Patch compute_num_split so _select_mhc_split_key_token_sizes works
-        _patch_compute_split(n_sms=188, monkeypatch=monkeypatch)
-
-        called_sizes: list[int] = []
-
-        class MockFn:
-            device = torch.device("cpu")
-
-        class MockLayer:
-            hidden_size = 4096
-            hc_mult = 16
-            hc_attn_fn = hc_attn_scale = hc_attn_base = MockFn()
-            hc_ffn_fn = hc_ffn_scale = hc_ffn_base = MockFn()
-
-            def hc_pre(self, residual_slice, fn, scale, base):
-                called_sizes.append(residual_slice.shape[0])
-                return (None, None, None)
-
-            def hc_post(self, layer_input, residual_slice, post_mix, comb_mix):
-                pass
-
-        MockLayer.__name__ = MockLayer.__qualname__ = "DeepseekV4DecoderLayer"
-
-        general_sizes = [1, 2, 4, 8, 16, 32, 64, 128]
-        _warmup_layer_mhc(MockLayer(), general_sizes)
-
-        # Compute expected union: general | split-key (K=65536, 188 SMs)
-        split_key_sizes = _select_mhc_split_key_token_sizes(
-            max_tokens=max(general_sizes), k_size=65536
-        )
-        expected_union = sorted(set(general_sizes) | set(split_key_sizes))
-
-        # hc_pre is called once per size per (attn + ffn) = 2x per size.
-        # Deduplicate to check unique sizes covered.
-        unique_called = sorted(set(called_sizes))
-        assert unique_called == expected_union, (
-            f"_warmup_layer_mhc called sizes: {unique_called}\n"
-            f"Expected union: {expected_union}\n"
-            f"General: {general_sizes}\n"
-            f"Split-key: {split_key_sizes}"
-        )
-
-    def test_calls_both_attn_and_ffn_for_each_size(self, monkeypatch) -> None:
-        from vllm.model_executor.warmup.deepseek_v4_mhc_warmup import (
-            _warmup_layer_mhc,
-        )
-
-        _patch_compute_split(n_sms=188, monkeypatch=monkeypatch)
-
-        call_log: list[str] = []
-
-        class MockFn:
-            device = torch.device("cpu")
-
-        class MockLayer:
-            hidden_size = 4096
-            hc_mult = 16
-            hc_attn_fn = hc_attn_scale = hc_attn_base = MockFn()
-            hc_ffn_fn = hc_ffn_scale = hc_ffn_base = MockFn()
-
-            def hc_pre(self, residual_slice, fn, scale, base):
-                size = residual_slice.shape[0]
-                call_log.append(f"hc_pre(size={size})")
-                return (None, None, None)
-
-            def hc_post(self, layer_input, residual_slice, post_mix, comb_mix):
-                call_log.append("hc_post")
-
-        MockLayer.__name__ = MockLayer.__qualname__ = "DeepseekV4DecoderLayer"
-
-        _patch_compute_split(n_sms=188, monkeypatch=monkeypatch)
-        _warmup_layer_mhc(MockLayer(), [1, 2])
-
-        # For each size: 2 calls (attn + ffn) * hc_pre+post = 4 log entries
-        # For 2 general sizes + any split-key reps bounded by max_tokens=2
-        # (which should be zero split-key reps since max_tokens=2 < 65)
-        # So 2 sizes * (hc_pre_attn, hc_post, hc_pre_ffn, hc_post) = 8 entries
-        assert len(call_log) == 8, (
-            f"Expected 8 log entries for 2 general sizes, got {len(call_log)}: "
-            f"{call_log}"
-        )
-
-
-# ── Model-type gate ───────────────────────────────────────────────────────
-
-
-class TestDeepseekV4ModelGate:
-    def test_returns_early_for_non_dsv4_model_type(self) -> None:
-        from vllm.model_executor.warmup.deepseek_v4_mhc_warmup import (
-            deepseek_v4_mhc_warmup,
-        )
-
-        class OtherModel:
-            config = SimpleNamespace(model_type="llama")
-
-            def modules(self):
-                return iter([])
-
-        deepseek_v4_mhc_warmup(OtherModel(), max_tokens=1024)
-
-
-# ── Orchestration ─────────────────────────────────────────────────────────
-
-
-class TestDeepseekV4MhcWarmupOrchestration:
-    """All three internal stages called in order.  TileLang dependencies
-    monkeypatched; mock modules use CUDA-like device attributes."""
-
-    def test_all_three_stages_called_in_order(self, monkeypatch) -> None:
-        from vllm.model_executor.warmup.deepseek_v4_mhc_warmup import (
-            deepseek_v4_mhc_warmup,
-        )
-
-        calls: list[tuple] = []
-        _fake_cuda = SimpleNamespace(device=torch.device("cuda"))
-
-        def _record_layer(layer, token_sizes):
-            calls.append(("_warmup_layer_mhc", token_sizes))
-
-        def _record_broadcast(model_arg, token_sizes):
-            calls.append(("_warmup_broadcast_mhc", model_arg, token_sizes))
-
-        def _record_head(model_arg, token_sizes):
-            calls.append(("_warmup_hc_head", token_sizes))
-
-        class MockLayer:
-            hc_pre = hc_post = lambda *a: None
-            hc_attn_fn = hc_attn_scale = hc_attn_base = _fake_cuda
-            hc_ffn_fn = hc_ffn_scale = hc_ffn_base = _fake_cuda
-            hidden_size = 4096
-            hc_mult = 16
-
-        MockLayer.__name__ = MockLayer.__qualname__ = "DeepseekV4DecoderLayer"
-
-        class MockDsModel:
-            hc_head_fn = hc_head_scale = hc_head_base = _fake_cuda
-            config = SimpleNamespace(hidden_size=4096)
-            hc_mult = 16
-            hc_eps = rms_norm_eps = 1e-6
-
-        MockDsModel.__name__ = MockDsModel.__qualname__ = "DeepseekV4Model"
-
-        class MockModel:
-            config = SimpleNamespace(model_type="deepseek_v4")
-
-            def modules(self):
-                yield self
-                yield MockDsModel()
-                yield MockLayer()
-
-        monkeypatch.setattr(
-            "vllm.model_executor.warmup.deepseek_v4_mhc_warmup._warmup_layer_mhc",
-            _record_layer,
-        )
-        monkeypatch.setattr(
-            "vllm.model_executor.warmup.deepseek_v4_mhc_warmup._warmup_broadcast_mhc",
-            _record_broadcast,
-        )
-        monkeypatch.setattr(
-            "vllm.model_executor.warmup.deepseek_v4_mhc_warmup._warmup_hc_head",
-            _record_head,
-        )
-        monkeypatch.setattr(
-            "vllm.model_executor.warmup.deepseek_v4_mhc_warmup.torch.accelerator.synchronize",
-            lambda: calls.append(("synchronize",)),
-        )
-        monkeypatch.setattr(
-            "vllm.model_executor.warmup.deepseek_v4_mhc_warmup.logger.info",
-            lambda *a, **kw: None,
-        )
-
-        deepseek_v4_mhc_warmup(MockModel(), max_tokens=1024)
-
-        assert len(calls) == 4, f"Expected 4 calls, got {len(calls)}: {calls}"
-
-        s1_name, s1_sizes = calls[0]
-        assert s1_name == "_warmup_layer_mhc"
-        assert isinstance(s1_sizes, list) and len(s1_sizes) > 0
-        assert 1 in s1_sizes
-
-        s2_name, s2_model, s2_sizes = calls[1]
-        assert s2_name == "_warmup_broadcast_mhc"
-        assert s2_sizes == s1_sizes
-        assert s2_model.__class__.__name__ == "MockModel"
-
-        s3_name, s3_sizes = calls[2]
-        assert s3_name == "_warmup_hc_head"
-        assert s3_sizes == s1_sizes
-
-        assert calls[3] == ("synchronize",)
-
-
-# ── No sys.modules leakage ───────────────────────────────────────────────
-
-
-class TestNoSysModulesLeakage:
-    """Verify that ``_select_mhc_split_key_token_sizes`` does not leave
-    ``tilelang_kernels`` permanently cached in ``sys.modules`` after the
-    stub is restored by monkeypatch cleanup."""
-
-    MODULE_PATH = "vllm.model_executor.kernels.mhc.tilelang_kernels"
-
-    def test_clean_modules_after_monkeypatch_cleanup(self, monkeypatch) -> None:
-        # Ensure module is not already in sys.modules
-        monkeypatch.delitem(sys.modules, self.MODULE_PATH, raising=False)
-
-        # Patch and invoke
-        _patch_compute_split(n_sms=188, monkeypatch=monkeypatch)
-        reps = _select_mhc_split_key_token_sizes(max_tokens=8192, k_size=4096)
-        assert len(reps) == 16
-
-        # Monkeypatch cleanup restores original state: module removed if absent
-        monkeypatch.undo()
-        # After undo, the module should not be present (it wasn't before)
-        assert self.MODULE_PATH not in sys.modules, (
-            f"{self.MODULE_PATH} leaked into sys.modules"
-        )

tests/v1/sample/test_head_dtype.py

@@ -14,7 +14,6 @@ import torch
 from vllm import LLM, SamplingParams
 from vllm.model_executor.layers.logits_processor import LogitsProcessor
 from vllm.model_executor.layers.vocab_parallel_embedding import (
-    ParallelLMHead,
     UnquantizedEmbeddingMethod,
 )
 
@@ -29,7 +28,6 @@ class _FakeLmHead:
         self.weight = weight
         self.quant_method = object() if quantized else UnquantizedEmbeddingMethod()
         self.shard_indices = shard_indices
-        self.tp_size = 1
 
 
 def _build_processor(vocab_size: int) -> LogitsProcessor:
@@ -137,59 +135,11 @@ def test_fp32_head_rejects_quantized_lm_head(default_vllm_config):
         lp._get_logits(torch.randn(4, 16, dtype=torch.bfloat16), lm_head, None)
 
 
-def test_replicated_lm_head_skips_tp_communication_and_preserves_processing(
-    default_vllm_config,
-):
-    from unittest import mock
-
-    vocab_size, hidden_size = 12, 8
-    soft_cap, scale = 2.0, 0.5
-    lp = LogitsProcessor(
-        vocab_size,
-        soft_cap=soft_cap,
-        scale=scale,
-    )
-    lp.head_dtype = torch.float32
-
-    hidden_states = torch.randn(4, hidden_size, dtype=torch.bfloat16)
-    weight = torch.randn(vocab_size, hidden_size, dtype=torch.bfloat16)
-    world_size_getter = (
-        "vllm.model_executor.layers.vocab_parallel_embedding."
-        "get_tensor_model_parallel_world_size"
-    )
-    with mock.patch(world_size_getter, return_value=2):
-        lm_head = ParallelLMHead(
-            vocab_size,
-            hidden_size,
-            params_dtype=torch.bfloat16,
-            disable_tp=True,
-        )
-    lm_head.weight_loader(lm_head.weight, weight)
-    assert lm_head.tp_size == 1
-
-    with mock.patch.object(lp, "_gather_logits") as gather_mock:
-        logits = lp(lm_head, hidden_states)
-
-    gather_mock.assert_not_called()
-
-    expected = torch.nn.functional.linear(hidden_states.float(), weight.float())
-    expected = torch.tanh(expected / soft_cap) * soft_cap * scale
-    torch.testing.assert_close(logits, expected)
-
-    all_gather_path = (
-        "vllm.model_executor.layers.logits_processor.tensor_model_parallel_all_gather"
-    )
-    with mock.patch(all_gather_path) as all_gather:
-        top = lp.get_top_tokens(lm_head, hidden_states)
-
-    all_gather.assert_not_called()
-    assert torch.equal(top, expected.argmax(dim=-1))
-
-
 def test_get_top_tokens_honors_head_dtype(default_vllm_config):
     # The spec-decode local-argmax path (get_top_tokens) must run the lm_head
     # in head_dtype too, not just _get_logits.
     import types
+    from unittest import mock
 
     vocab_size, hidden_size = 64, 16
     lp = _build_processor(vocab_size)
@@ -204,7 +154,13 @@ def test_get_top_tokens_honors_head_dtype(default_vllm_config):
         ),
     )
 
-    top = lp.get_top_tokens(lm_head, hidden_states, None)
+    with mock.patch(
+        "vllm.model_executor.layers.logits_processor."
+        "get_tensor_model_parallel_world_size",
+        return_value=1,
+    ):
+        top = lp.get_top_tokens(lm_head, hidden_states, None)
+
     expected = torch.nn.functional.linear(hidden_states.float(), weight.float()).argmax(
         dim=-1
     )

vllm/compilation/passes/fx_utils.py

@@ -14,10 +14,6 @@ def is_func(node: fx.Node, target: Target) -> bool:
     return bool(node.op == "call_function" and node.target == target)
 
 
-def is_auto_func(node: fx.Node, op: OpOverload) -> bool:
-    return is_func(node, auto_functionalized) and node.args[0] == op
-
-
 # Returns the first auto_functionalized node with the given op (if it exists)
 def find_auto_fn_maybe(nodes: Iterable[fx.Node], op: OpOverload) -> fx.Node | None:
     for node in nodes:
@@ -42,13 +38,6 @@ def find_getitem_maybe(node: fx.Node, idx: int) -> fx.Node | None:
     return None
 
 
-# Returns the getitem node that extracts the idx-th element from node
-def find_getitem(node: fx.Node, idx: int) -> fx.Node:
-    ret = find_getitem_maybe(node, idx)
-    assert ret is not None, f"Could not find getitem {idx} in node {node}"
-    return ret
-
-
 # An auto-functionalization-aware utility for finding nodes with a specific op
 # Also handles op overload packets and finds all overloads
 def find_op_nodes(
@@ -67,11 +56,3 @@ def find_op_nodes(
     for n in graph.find_nodes(op="call_function", target=auto_functionalized):
         if n.args[0] == op:
             yield n
-
-
-# Asserts that the node only has one user and returns it
-# Even if a node has only 1 user, it might share storage with another node,
-# which might need to be taken into account.
-def get_only_user(node: fx.Node) -> fx.Node:
-    assert len(node.users) == 1
-    return next(iter(node.users))

vllm/compilation/passes/vllm_inductor_pass.py

@@ -333,12 +333,3 @@ class VllmFusionPatternMatcherPass(VllmPatternMatcherPass):
     def __call__(self, graph: torch.fx.Graph) -> None:
         self.matched_count = self.pm_pass.apply(graph)
         VllmPatternMatcherPass.match_table[self.pass_name] += self.matched_count
-
-
-class PrinterInductorPass(VllmInductorPass):
-    def __init__(self, name: str, config: VllmConfig) -> None:
-        super().__init__(config)
-        self.name = name
-
-    def __call__(self, graph: torch.fx.Graph) -> None:
-        self.dump_graph(graph, self.name)

vllm/config/utils.py

@@ -415,31 +415,6 @@ class Range:
         return self.__str__()
 
 
-def handle_deprecated(
-    config: ConfigT,
-    old_name: str,
-    new_name_or_names: str | list[str],
-    removal_version: str,
-) -> None:
-    old_val = getattr(config, old_name)
-    if old_val is None:
-        return
-
-    if isinstance(new_name_or_names, str):
-        new_names = [new_name_or_names]
-    else:
-        new_names = new_name_or_names
-
-    msg = (
-        f"{old_name} is deprecated and will be removed in {removal_version}. "
-        f"Use {', '.join(new_names)} instead."
-    )
-    logger.warning(msg)
-
-    for new_name in new_names:
-        setattr(config, new_name, old_val)
-
-
 def get_from_deprecated_env_if_set(
     env_name: str,
     removal_version: str,

vllm/distributed/parallel_state.py

@@ -2081,9 +2081,6 @@ def model_parallel_is_initialized():
     return _TP is not None and _PP is not None
 
 
-_TP_STATE_PATCHED = False
-
-
 def get_tensor_model_parallel_world_size() -> int:
     """Return world size for the tensor model parallel group."""
     return get_tp_group().world_size

vllm/entrypoints/chat_utils.py

@@ -2015,15 +2015,6 @@ async def parse_chat_messages_async(
     return conversation, mm_data, mm_uuids
 
 
-def get_history_tool_calls_cnt(conversation: list[ConversationMessage]):
-    idx = 0
-    for msg in conversation:
-        if msg["role"] == "assistant":
-            tool_calls = msg.get("tool_calls")
-            idx += len(list(tool_calls)) if tool_calls is not None else 0  # noqa
-    return idx
-
-
 _KIMI_MODEL_TYPES = ("kimi_k2", "kimi_k25", "kimi_k3")
 
 

vllm/model_executor/layers/activation.py

@@ -8,14 +8,8 @@ import torch
 import torch.nn as nn
 import torch.nn.functional as F
 
-from vllm.distributed import (
-    divide,
-    get_tensor_model_parallel_rank,
-    get_tensor_model_parallel_world_size,
-)
 from vllm.logger import init_logger
 from vllm.model_executor.custom_op import CustomOp
-from vllm.model_executor.utils import set_weight_attrs
 from vllm.platforms import CpuArchEnum, current_platform
 from vllm.triton_utils import tl, triton
 from vllm.utils.collection_utils import LazyDict
@@ -788,48 +782,6 @@ class XIELU(CustomOp):
         return self.forward_native(input)
 
 
-class ScaledActivation(nn.Module):
-    """An activation function with post-scale parameters.
-
-    This is used for some quantization methods like AWQ.
-    """
-
-    def __init__(
-        self,
-        act_module: nn.Module,
-        intermediate_size: int,
-        input_is_parallel: bool = True,
-        params_dtype: torch.dtype | None = None,
-    ):
-        super().__init__()
-        self.act = act_module
-        self.input_is_parallel = input_is_parallel
-        if input_is_parallel:
-            tp_size = get_tensor_model_parallel_world_size()
-            intermediate_size_per_partition = divide(intermediate_size, tp_size)
-        else:
-            intermediate_size_per_partition = intermediate_size
-        if params_dtype is None:
-            params_dtype = torch.get_default_dtype()
-        self.scales = nn.Parameter(
-            torch.empty(intermediate_size_per_partition, dtype=params_dtype)
-        )
-        set_weight_attrs(self.scales, {"weight_loader": self.weight_loader})
-
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        return self.act(x) / self.scales
-
-    def weight_loader(self, param: nn.Parameter, loaded_weight: torch.Tensor):
-        param_data = param.data
-        if self.input_is_parallel:
-            tp_rank = get_tensor_model_parallel_rank()
-            shard_size = param_data.shape[0]
-            start_idx = tp_rank * shard_size
-            loaded_weight = loaded_weight.narrow(0, start_idx, shard_size)
-        assert param_data.shape == loaded_weight.shape
-        param_data.copy_(loaded_weight)
-
-
 _ACTIVATION_REGISTRY = LazyDict(
     {
         "gelu": lambda: GELU(),

vllm/model_executor/layers/conv.py

@@ -155,58 +155,6 @@ class Conv2dLayer(ConvLayerBase):
         return self._forward_conv(x)
 
 
-class CausalConv2dLayer(Conv2dLayer):
-    """
-    A causal version of nn.Conv2d where each location in the 2D matrix would
-    have no access to locations on its right or down
-    All arguments are the same as nn.Conv2d except padding which should be
-    set as None
-    """
-
-    def __init__(
-        self,
-        in_channels: int,
-        out_channels: int,
-        kernel_size: int,
-        stride: int,
-        padding: int = 0,
-        dilation: int = 1,
-        groups: int = 1,
-        bias: bool = True,
-        padding_mode: str = "zeros",
-        *,
-        params_dtype: torch.dtype | None = None,
-    ) -> None:
-        if padding is not None:
-            raise ValueError(
-                "Argument padding should be set to None for CausalConv2dLayer."
-            )
-        self._left_padding: int = kernel_size - 1
-        self._right_padding: int = stride - 1
-        padding = 0
-
-        super().__init__(
-            in_channels,
-            out_channels,
-            kernel_size,
-            stride,
-            padding,
-            dilation,
-            groups,
-            bias,
-            padding_mode,
-            params_dtype=params_dtype,
-        )
-
-    def forward(
-        self,
-        x: torch.Tensor,
-    ) -> torch.Tensor:
-        x = F.pad(x, pad=(self._left_padding, self._right_padding, 0, 0))
-        x = super().forward(x)
-        return x
-
-
 # --8<-- [start:conv3d]
 @CustomOp.register("conv3d")
 class Conv3dLayer(ConvLayerBase):

vllm/model_executor/layers/logits_processor.py

@@ -7,6 +7,7 @@ import torch.nn.functional as F
 
 from vllm.config import get_current_vllm_config
 from vllm.distributed import (
+    get_tensor_model_parallel_world_size,
     tensor_model_parallel_all_gather,
     tensor_model_parallel_gather,
 )
@@ -144,8 +145,7 @@ class LogitsProcessor(PluggableLayer):
         logits = self._apply_head(lm_head, hidden_states, embedding_bias)
 
         # Gather logits for TP
-        if lm_head.tp_size > 1:
-            logits = self._gather_logits(logits)
+        logits = self._gather_logits(logits)
 
         # Remove paddings in vocab (if any).
         if logits is not None:
@@ -169,7 +169,7 @@ class LogitsProcessor(PluggableLayer):
                 "The local argmax reduction optimization is not supported for "
                 "non-positive logit scaling factors."
             )
-        tp_size = lm_head.tp_size
+        tp_size = get_tensor_model_parallel_world_size()
 
         logits = self._apply_head(lm_head, hidden_states, embedding_bias)
         if self.soft_cap is not None:

vllm/model_executor/layers/quantization/fp8.py

@@ -4,7 +4,6 @@
 from typing import TYPE_CHECKING, Any
 
 import torch
-from torch.utils._python_dispatch import TorchDispatchMode
 
 import vllm.envs as envs
 import vllm.model_executor.layers.fused_moe.modular_kernel as mk
@@ -234,26 +233,6 @@ class Fp8Config(QuantizationConfig):
         return cache_scale_mapper | QuantizationConfig.get_cache_scale_mapper()
 
 
-class CopyNumelCounter(TorchDispatchMode):
-    """
-    Tracks total number of elements modified with `copy_`. Useful for keeping
-    track of weight loading where underlying weights can be arbitrarily
-    transformed (such as with `narrow`) before calling copy.
-    """
-
-    def __init__(self):
-        super().__init__()
-        self.copied_numel = 0
-
-    def __torch_dispatch__(self, func, types, args=(), kwargs=None):
-        if kwargs is None:
-            kwargs = {}
-        out = func(*args, **kwargs)
-        if func == torch.ops.aten.copy_.default:
-            self.copied_numel += args[0].numel()
-        return out
-
-
 class Fp8LinearMethod(LinearMethodBase):
     """Linear method for FP8.
     Supports loading FP8 checkpoints with static weight scale and

vllm/model_executor/layers/utils.py

@@ -25,12 +25,6 @@ MOE_LAYER_ROUTER_GATE_SUFFIXES = {
 }
 
 
-def is_layer_moe_router_gate(prefix: str) -> bool:
-    if not prefix:
-        return False
-    return prefix.rsplit(".", 1)[-1] in MOE_LAYER_ROUTER_GATE_SUFFIXES
-
-
 def get_token_bin_counts_and_mask(
     tokens: torch.Tensor,
     vocab_size: int,

vllm/model_executor/layers/vocab_parallel_embedding.py

@@ -232,7 +232,6 @@ class VocabParallelEmbedding(PluggableLayer):
         padding_size: padding size for the vocabulary.
         quant_config: quant config for the layer
         prefix: full name of the layer in the state dict
-        disable_tp: If true, tensor parallelism will be disabled for this layer.
     """  # noqa: E501
 
     # --8<-- [end:vocab_parallel_embedding]
@@ -246,19 +245,12 @@ class VocabParallelEmbedding(PluggableLayer):
         padding_size: int = DEFAULT_VOCAB_PADDING_SIZE,
         quant_config: QuantizationConfig | None = None,
         prefix: str = "",
-        *,
-        disable_tp: bool = False,
     ):
         super().__init__()
 
         # Keep the input dimensions.
-        self.disable_tp = disable_tp
-        if disable_tp:
-            tp_rank, self.tp_size = 0, 1
-        else:
-            tp_rank = get_tensor_model_parallel_rank()
-            self.tp_size = get_tensor_model_parallel_world_size()
-        self.tp_rank = tp_rank
+        tp_rank = get_tensor_model_parallel_rank()
+        self.tp_size = get_tensor_model_parallel_world_size()
         self.num_embeddings = num_embeddings
         self.padding_size = padding_size
         self.org_vocab_size = org_num_embeddings or num_embeddings
@@ -331,13 +323,6 @@ class VocabParallelEmbedding(PluggableLayer):
             params_dtype=params_dtype,
             weight_loader=self.weight_loader,
         )
-        self.update_param_tp_status()
-
-    def update_param_tp_status(self):
-        for param in self.parameters():
-            if isinstance(param, BasevLLMParameter):
-                param.tp_rank = self.tp_rank
-                param.tp_size = self.tp_size
 
     @classmethod
     def _get_indices(
@@ -502,9 +487,9 @@ class VocabParallelEmbedding(PluggableLayer):
         # Mask the output embedding.
         if self.tp_size > 1:
             output_parallel.masked_fill_(input_mask.unsqueeze(-1), 0)
-            # Reduce across all the model parallel GPUs.
-            return tensor_model_parallel_all_reduce(output_parallel)
-        return output_parallel
+        # Reduce across all the model parallel GPUs.
+        output = tensor_model_parallel_all_reduce(output_parallel)
+        return output
 
     def extra_repr(self) -> str:
         s = f"num_embeddings={self.num_embeddings_per_partition}"
@@ -531,7 +516,6 @@ class ParallelLMHead(VocabParallelEmbedding):
         params_dtype: type of the parameters.
         org_num_embeddings: original vocabulary size (without LoRA).
         padding_size: padding size for the vocabulary.
-        disable_tp: If true, tensor parallelism will be disabled for this layer.
     """
 
     # --8<-- [end:parallel_lm_head]
@@ -546,8 +530,6 @@ class ParallelLMHead(VocabParallelEmbedding):
         padding_size: int = DEFAULT_VOCAB_PADDING_SIZE,
         quant_config: QuantizationConfig | None = None,
         prefix: str = "",
-        *,
-        disable_tp: bool = False,
     ):
         super().__init__(
             num_embeddings,
@@ -557,7 +539,6 @@ class ParallelLMHead(VocabParallelEmbedding):
             padding_size,
             quant_config,
             prefix,
-            disable_tp=disable_tp,
         )
         self.quant_config = quant_config
         if bias:

vllm/model_executor/models/glm_ocr.py

@@ -242,10 +242,6 @@ class GlmOcrVisionBlock(Glm4vVisionBlock):
         )
 
 
-class GlmOcrVisionPatchEmbed(Glm4vVisionPatchEmbed):
-    pass
-
-
 class GlmOcrPatchMerger(Glm4vPatchMerger):
     pass
 

vllm/model_executor/models/idefics3.py

@@ -103,41 +103,6 @@ class Idefics3ProcessingInfo(BaseProcessingInfo):
     def get_supported_mm_limits(self) -> Mapping[str, int | None]:
         return {"image": None}
 
-    def _resize_output_size(
-        self,
-        *,
-        height: int,
-        width: int,
-        max_len: int | None = None,
-        min_len: int = 1,
-        max_size: int | None = None,
-    ) -> tuple[int, int]:
-        # Set default value for max_len if not provided
-        max_len = max(height, width) if max_len is None else max_len
-        aspect_ratio = width / height
-
-        # Handle the maximum size constraint
-        if max_size is not None:
-            max_len = min(max_len, max_size)
-
-        # Adjust dimensions according to the aspect ratio
-        if width >= height:
-            width = max_len
-            height = int(width / aspect_ratio)
-        else:
-            height = max_len
-            width = int(height * aspect_ratio)
-
-        # Ensure both width and height are even (if needed)
-        height += height % 2
-        width += width % 2
-
-        # Ensure dimensions are not smaller than the minimum length
-        height = max(height, min_len)
-        width = max(width, min_len)
-
-        return height, width
-
     def _get_image_feature_grid_size(
         self,
         *,

vllm/model_executor/models/llava_onevision2.py

@@ -1531,13 +1531,6 @@ class LlavaOnevision2MultiModalDataParser(MultiModalDataParser):
 class LlavaOnevision2MultiModalProcessor(
     BaseMultiModalProcessor[LlavaOnevision2ProcessingInfo]
 ):
-    def _get_data_parser(self) -> MultiModalDataParser:
-        # Retained for symmetry; vLLM actually fetches the parser via
-        # info.get_data_parser() (see ProcessingInfo override above).
-        return LlavaOnevision2MultiModalDataParser(
-            self.info.get_hf_config().vision_config.spatial_merge_size
-        )
-
     def _call_hf_processor(
         self,
         prompt: str,

vllm/model_executor/models/nemotron_parse.py

@@ -77,24 +77,6 @@ class BartScaledWordEmbedding(VocabParallelEmbedding):
         return super().forward(input_ids) * self.embed_scale
 
 
-class BartParallelLMHead(ParallelLMHead):
-    """
-    This module overrides ParallelLMHead's
-    forward by dividing by embeddings scale,
-    yielding effectively the inverse of
-    BartScaledWordEmbedding
-    """
-
-    def __init__(
-        self, num_embeddings: int, embedding_dim: int, embed_scale: float = 1.0
-    ):
-        super().__init__(num_embeddings, embedding_dim)
-        self.embed_scale = embed_scale
-
-    def forward(self, input_ids: torch.Tensor) -> torch.Tensor:
-        return super().forward(input_ids) / self.embed_scale
-
-
 class BartDecoderLayer(nn.Module):
     def __init__(
         self,

vllm/model_executor/models/phi4mm.py

@@ -521,32 +521,6 @@ class Phi4MMAudioEmbeddingInputs(TensorSchema):
 Phi4MMAudioInputs: TypeAlias = Phi4MMAudioFeatureInputs | Phi4MMAudioEmbeddingInputs
 
 
-def cat_with_pad(tensors, dim, padding_value=0):
-    """
-    cat along dim, while pad to max for all other dims
-    """
-    ndim = tensors[0].dim()
-    assert all(t.dim() == ndim for t in tensors[1:]), (
-        "All tensors must have the same number of dimensions"
-    )
-
-    out_size = [max(t.shape[i] for t in tensors) for i in range(ndim)]
-    out_size[dim] = sum(t.shape[dim] for t in tensors)
-    output = tensors[0].new_full(out_size, padding_value)
-
-    index = 0
-    for t in tensors:
-        # Create a slice list where every dimension except dim is full slice
-        slices = [slice(0, t.shape[d]) for d in range(ndim)]
-        # Update only the concat dimension slice
-        slices[dim] = slice(index, index + t.shape[dim])
-
-        output[slices] = t
-        index += t.shape[dim]
-
-    return output
-
-
 def stack_with_pad(
     tensors: torch.Tensor | list[torch.Tensor],
     padding_value: int | float = 0,

vllm/model_executor/models/phi4mm_utils.py

@@ -1595,14 +1595,6 @@ class AttModule(nn.Module):
         return x, memory, pos_emb, att_mask
 
 
-class AttBlock(BlockBase, AttModule):
-    """Attention Block module to support both Attention and Block module."""
-
-    def memory_dims(self, max_len: bool = False) -> tuple[int, int]:
-        """memory dimensions"""
-        return (1, self.input_size)
-
-
 def masked_softmax(
     scores: Tensor,
     mask: Tensor | None,

vllm/model_executor/models/qwen3_dspark.py

@@ -24,6 +24,7 @@ from vllm.logger import init_logger
 from vllm.model_executor.layers.logits_processor import LogitsProcessor
 from vllm.model_executor.layers.vocab_parallel_embedding import (
     ParallelLMHead,
+    VocabParallelEmbedding,
 )
 
 from .qwen3_dflash import DFlashQwen3ForCausalLM, DFlashQwen3Model
@@ -39,10 +40,6 @@ class DSparkMarkovHead(nn.Module):
     ``vocab_size``); ``markov_w2`` projects it to a draft-vocab bias
     (``draft_vocab_size``) added to the base draft logits. The two sizes
     coincide for full-vocab drafts.
-
-    Both weights are replicated because the head runs sequentially for every
-    draft position. Sharding them would add an all-reduce and a full-vocab
-    gather to each position.
     """
 
     def __init__(
@@ -53,24 +50,19 @@ class DSparkMarkovHead(nn.Module):
         prefix: str,
     ) -> None:
         super().__init__()
-        self.markov_w1 = nn.Embedding(vocab_size, markov_rank)
+        # TODO(ben): profile for which (if any) it makes sense to replicate or TP-shard
+        self.markov_w1 = VocabParallelEmbedding(
+            vocab_size, markov_rank, prefix=maybe_prefix(prefix, "markov_w1")
+        )
         self.markov_w2 = ParallelLMHead(
-            draft_vocab_size,
-            markov_rank,
-            bias=False,
-            prefix=maybe_prefix(prefix, "markov_w2"),
-            disable_tp=True,
+            draft_vocab_size, markov_rank, prefix=maybe_prefix(prefix, "markov_w2")
         )
 
     def embed(self, token_ids: torch.Tensor) -> torch.Tensor:
         """r-dim Markov embedding of ``token_ids`` ([B] -> [B, r])."""
         return self.markov_w1(token_ids)
 
-    def bias(
-        self,
-        markov_embed: torch.Tensor,
-        logits_processor: LogitsProcessor,
-    ) -> torch.Tensor:
+    def bias(self, markov_embed: torch.Tensor, logits_processor) -> torch.Tensor:
         """Vocab-size transition bias from a Markov embedding ([B, r] -> [B, V])."""
         return logits_processor(self.markov_w2, markov_embed)
 

vllm/model_executor/models/skyworkr1v.py

@@ -7,7 +7,7 @@
 # Copyright (c) 2025 Skywork
 # Licensed under The MIT License [see LICENSE for details]
 # --------------------------------------------------------
-from collections.abc import Iterable, Mapping
+from collections.abc import Iterable
 from typing import Annotated, Literal, TypeAlias
 
 import torch
@@ -15,8 +15,6 @@ import torch.nn as nn
 from transformers import PretrainedConfig
 
 from vllm.config import VllmConfig
-from vllm.config.multimodal import BaseDummyOptions
-from vllm.inputs import MultiModalDataDict
 from vllm.model_executor.layers.linear import ReplicatedLinear
 from vllm.model_executor.layers.quantization import QuantizationConfig
 from vllm.model_executor.layers.quantization.auto_awq import AutoAWQConfig
@@ -24,7 +22,6 @@ from vllm.model_executor.models.intern_vit import (
     InternVisionModel,
 )
 from vllm.multimodal import MULTIMODAL_REGISTRY
-from vllm.multimodal.processing import BaseDummyInputsBuilder
 from vllm.sequence import IntermediateTensors
 from vllm.transformers_utils.processors.internvl import (
     InternVLImageProcessor,
@@ -117,33 +114,6 @@ class SkyworkR1VProcessingInfo(BaseInternVLProcessingInfo):
         )
 
 
-class SkyworkR1VDummyInputsBuilder(BaseDummyInputsBuilder[SkyworkR1VProcessingInfo]):
-    def get_dummy_text(self, mm_counts: Mapping[str, int]) -> str:
-        num_images = mm_counts.get("image", 0)
-
-        return "<image>" * num_images
-
-    def get_dummy_mm_data(
-        self,
-        seq_len: int,
-        mm_counts: Mapping[str, int],
-        mm_options: Mapping[str, BaseDummyOptions],
-    ) -> MultiModalDataDict:
-        target_width, target_height = self.info.get_image_size_with_most_features()
-        num_images = mm_counts.get("image", 0)
-
-        image_overrides = mm_options.get("image")
-
-        return {
-            "image": self._get_dummy_images(
-                width=target_width,
-                height=target_height,
-                num_images=num_images,
-                overrides=image_overrides,
-            )
-        }
-
-
 @MULTIMODAL_REGISTRY.register_processor(
     BaseInternVLMultiModalProcessor,
     info=SkyworkR1VProcessingInfo,

vllm/model_executor/models/utils.py

@@ -27,7 +27,6 @@ from vllm.multimodal import NestedTensors
 from vllm.sequence import IntermediateTensors
 from vllm.utils.math_utils import cdiv
 from vllm.utils.torch_utils import (
-    async_tensor_h2d,
     direct_register_custom_op,
 )
 
@@ -673,17 +672,6 @@ def _merge_multimodal_embeddings(
     return inputs_embeds
 
 
-def isin_list(
-    elements: torch.Tensor,
-    test_elements_list: list[int],
-) -> torch.Tensor:
-    test_elements = async_tensor_h2d(
-        test_elements_list, dtype=torch.int64, device=elements.device
-    )
-
-    return torch.isin(elements, test_elements)
-
-
 class StageMissingLayer(nn.Module):
     def __init__(self, stage_name: str, module: nn.Module | None = None) -> None:
         super().__init__()

vllm/model_executor/models/vision.py

@@ -580,40 +580,3 @@ def run_dp_sharded_mrope_vision_model(
         "Found unassigned embeddings"
     )
     return out_embeddings
-
-
-def get_llm_pos_ids_for_vision(
-    start_idx: int,
-    vision_idx: int,
-    spatial_merge_size: int,
-    t_index: list[int],
-    grid_hs: torch.Tensor,
-    grid_ws: torch.Tensor,
-) -> torch.Tensor:
-    llm_pos_ids_list = []
-    llm_grid_h = grid_hs[vision_idx] // spatial_merge_size
-    llm_grid_w = grid_ws[vision_idx] // spatial_merge_size
-    h_index = (
-        torch.arange(llm_grid_h)
-        .view(1, -1, 1)
-        .expand(len(t_index), -1, llm_grid_w)
-        .flatten()
-    )
-    w_index = (
-        torch.arange(llm_grid_w)
-        .view(1, 1, -1)
-        .expand(len(t_index), llm_grid_h, -1)
-        .flatten()
-    )
-    t_index_tensor = (
-        torch.Tensor(t_index)
-        .to(llm_grid_h.device)
-        .view(-1, 1)
-        .expand(-1, llm_grid_h * llm_grid_w)
-        .long()
-        .flatten()
-    )
-    _llm_pos_ids = torch.stack([t_index_tensor, h_index, w_index])
-    llm_pos_ids_list.append(_llm_pos_ids + start_idx)
-    llm_pos_ids = torch.cat(llm_pos_ids_list, dim=1)
-    return llm_pos_ids

vllm/model_executor/warmup/deepseek_v4_mhc_warmup.py

@@ -15,7 +15,6 @@ import torch
 
 from vllm.logger import init_logger
 from vllm.tracing import instrument
-from vllm.utils.math_utils import cdiv
 
 logger = init_logger(__name__)
 
@@ -62,39 +61,6 @@ def _select_mhc_warmup_token_sizes(
     return _normalize_token_sizes(candidates, max_tokens=max_auto_tokens)
 
 
-def _select_mhc_split_key_token_sizes(
-    *,
-    max_tokens: int,
-    k_size: int,
-) -> list[int]:
-    """Select one representative token count per distinct n_splits compile key.
-
-    The MHC TileLang kernels compute n_splits at runtime via
-    ``compute_num_split(block_k, k_size, cdiv(tokens, block_k))``.  Because
-    ``n_splits`` is a TileLang compile-time parameter, every distinct value
-    produces a separate JIT compilation artifact.  This function returns
-    exactly one token per reachable ``n_splits`` value, stopping early when
-    ``n_splits`` drops to 1 (all remaining grid sizes map to n_splits=1).
-    """
-    from vllm.model_executor.kernels.mhc.tilelang_kernels import compute_num_split
-
-    block_k = 64
-    max_grid = cdiv(max_tokens, block_k)
-    reps: list[int] = []
-    seen: set[int] = set()
-    for g in range(1, max_grid + 1):
-        t = (g - 1) * block_k + 1
-        if t > max_tokens:
-            break
-        ns = compute_num_split(block_k, k_size, g)
-        if ns not in seen:
-            reps.append(t)
-            seen.add(ns)
-            if ns == 1:
-                break
-    return reps
-
-
 def _find_first_mhc_layer(model: torch.nn.Module) -> torch.nn.Module | None:
     for module in model.modules():
         if module.__class__.__name__ != "DeepseekV4DecoderLayer":
@@ -136,16 +102,6 @@ def _warmup_layer_mhc(
     hidden_size = int(layer.hidden_size)
     hc_mult = int(layer.hc_mult)
     device = layer.hc_attn_fn.device
-
-    # Union general token sizes with split-key reps for the non-broadcast
-    # MHC kernel (k_size = hc_mult * hidden_size).
-    k_size = hc_mult * hidden_size
-    split_key_sizes = _select_mhc_split_key_token_sizes(
-        max_tokens=max_tokens, k_size=k_size
-    )
-    all_sizes = sorted(set(token_sizes) | set(split_key_sizes))
-    max_tokens = max(all_sizes)
-
     residual = torch.zeros(
         max_tokens,
         hc_mult,
@@ -154,19 +110,19 @@ def _warmup_layer_mhc(
         device=device,
     )
 
-    for size in all_sizes:
+    for size in token_sizes:
         residual_slice = residual[:size]
         for fn, scale, base in (
             (layer.hc_attn_fn, layer.hc_attn_scale, layer.hc_attn_base),
             (layer.hc_ffn_fn, layer.hc_ffn_scale, layer.hc_ffn_base),
         ):
-            layer_input, post_mix, res_mix = layer.hc_pre(
+            layer_input, post_mix, comb_mix = layer.hc_pre(
                 residual_slice,
                 fn,
                 scale,
                 base,
             )
-            layer.hc_post(layer_input, residual_slice, post_mix, res_mix)
+            layer.hc_post(layer_input, residual_slice, post_mix, comb_mix)
 
 
 def _warmup_hc_head(
@@ -205,86 +161,6 @@ def _warmup_hc_head(
         )
 
 
-def _warmup_broadcast_mhc(
-    model: torch.nn.Module,
-    token_sizes: list[int],
-) -> None:
-    """Warm up the first-layer broadcast MHC TileLang kernel.
-
-    The first ``DeepseekV4DecoderLayer`` uses
-    ``mhc_pre_broadcast_tilelang`` (2-D input, ``fn_broadcast`` weight)
-    instead of the 3-D ``mhc_pre_tilelang`` used by all subsequent layers.
-    ``fn_broadcast`` is set during ``finalize_mhc_broadcast_weights()`` and
-    only exists on the very first decoder layer.  No-op for models without a
-    broadcast-capable layer.
-
-    Unlike the generic per-layer MHC warmup (which covers power-of-two token
-    sizes), the broadcast kernel uses ``n_splits`` as a TileLang compile-time
-    parameter.  Different token counts can map to the same ``n_splits`` value;
-    this function selects one representative token per distinct compile key to
-    avoid redundant JIT compilations while covering every reachable key.
-    """
-    first_broadcast_layer = None
-    for module in model.modules():
-        if module.__class__.__name__ != "DeepseekV4DecoderLayer":
-            continue
-        fn_broadcast = getattr(module, "hc_attn_fn_broadcast", None)
-        if fn_broadcast is not None:
-            first_broadcast_layer = module
-            break
-    if first_broadcast_layer is None:
-        return
-
-    from vllm.model_executor.kernels.mhc.tilelang import mhc_pre_broadcast_tilelang
-
-    device = first_broadcast_layer.hc_attn_fn.device
-    if device.type != "cuda":
-        return
-
-    # Fail closed unless hc_attn_fn_broadcast is a tensor on the same CUDA
-    # device as hc_attn_fn.  A non-tensor, a CPU tensor, or a tensor on a
-    # different device index indicates a partially initialized or mismatched
-    # broadcast-weight setup that must not be warmed up.
-    fn_broadcast = first_broadcast_layer.hc_attn_fn_broadcast
-    if not isinstance(fn_broadcast, torch.Tensor):
-        return
-    if fn_broadcast.device.type != "cuda":
-        return
-    if fn_broadcast.device != device:
-        return
-
-    hidden_size = first_broadcast_layer.hidden_size
-    broadcast_token_sizes = _select_mhc_split_key_token_sizes(
-        max_tokens=max(token_sizes),
-        k_size=hidden_size,
-    )
-
-    x_2d = torch.zeros(
-        max(broadcast_token_sizes),
-        hidden_size,
-        dtype=torch.bfloat16,
-        device=device,
-    )
-    norm_weight = first_broadcast_layer.attn_norm.weight.data
-    norm_eps = first_broadcast_layer.attn_norm.variance_epsilon
-
-    for size in broadcast_token_sizes:
-        mhc_pre_broadcast_tilelang(
-            x_2d[:size],
-            first_broadcast_layer.hc_attn_fn,
-            first_broadcast_layer.hc_attn_scale,
-            first_broadcast_layer.hc_attn_base,
-            first_broadcast_layer.rms_norm_eps,
-            first_broadcast_layer.hc_eps,
-            first_broadcast_layer.hc_eps,
-            first_broadcast_layer.hc_post_alpha,
-            first_broadcast_layer.hc_sinkhorn_iters,
-            norm_weight=norm_weight,
-            norm_eps=norm_eps,
-            fn_broadcast=first_broadcast_layer.hc_attn_fn_broadcast,
-        )
-
-
 @instrument(span_name="DeepSeek V4 mHC warmup")
 def deepseek_v4_mhc_warmup(
     model: torch.nn.Module,
@@ -323,7 +199,6 @@ def deepseek_v4_mhc_warmup(
     )
     with torch.inference_mode():
         _warmup_layer_mhc(layer, token_sizes)
-        _warmup_broadcast_mhc(model, token_sizes)
         if deepseek_model is not None:
             _warmup_hc_head(deepseek_model, token_sizes)
         torch.accelerator.synchronize()

vllm/models/deepseek_v4/nvidia/model.py

@@ -865,43 +865,6 @@ class DeepseekV4DecoderLayer(nn.Module):
             requires_grad=False,
         )
 
-    def hc_pre(
-        self,
-        x: torch.Tensor,
-        hc_fn: torch.Tensor,
-        hc_scale: torch.Tensor,
-        hc_base: torch.Tensor,
-    ) -> tuple[torch.Tensor, torch.Tensor, torch.Tensor]:
-        """Run standalone MHC pre (non-broadcast 3D path).
-
-        Used by kernel warmup and model-level forward.  Passes fused RMSNorm
-        weight so the TileLang JIT compiles the production kernel path.
-        """
-        post_mix, res_mix, layer_input = mhc_pre_tilelang(
-            residual=x,
-            fn=hc_fn,
-            hc_scale=hc_scale,
-            hc_base=hc_base,
-            rms_eps=self.rms_norm_eps,
-            hc_pre_eps=self.hc_eps,
-            hc_sinkhorn_eps=self.hc_eps,
-            hc_post_mult_value=self.hc_post_alpha,
-            sinkhorn_repeat=self.hc_sinkhorn_iters,
-            norm_weight=self.attn_norm.weight.data,
-            norm_eps=self.attn_norm.variance_epsilon,
-        )
-        return layer_input, post_mix, res_mix
-
-    def hc_post(
-        self,
-        x: torch.Tensor,
-        residual: torch.Tensor,
-        post: torch.Tensor,
-        comb: torch.Tensor,
-    ) -> torch.Tensor:
-        """Run standalone MHC post.  Used by kernel warmup."""
-        return mhc_post_tilelang(x, residual, post, comb)
-
     def forward(
         self,
         x: torch.Tensor,

vllm/models/kimi_k3/nvidia/kda.py

@@ -164,8 +164,6 @@ def is_flashkda_supported(
     dtype: torch.dtype,
     lower_bound: float | None,
 ) -> bool:
-    if not current_platform.is_cuda():
-        return False
     capability = current_platform.get_device_capability()
     return (
         capability is not None

vllm/profiler/utils.py

@@ -12,15 +12,6 @@ from torch._C._profiler import _EventType, _ProfilerEvent, _TensorMetadata
 #
 
 
-def trim_string_front(string: str, width: int) -> str:
-    if len(string) > width:
-        offset = len(string) - width + 3
-        string = string[offset:]
-        if len(string) > 3:
-            string = "..." + string[3:]
-    return string
-
-
 def trim_string_back(string: str, width: int) -> str:
     if len(string) > width:
         offset = len(string) - width + 3

vllm/utils/cache.py

@@ -182,10 +182,6 @@ class LRUCache(cachetools.LRUCache[_K, _V]):
 
         self.popitem(remove_pinned=remove_pinned)
 
-    def _remove_old_if_needed(self) -> None:
-        while self.currsize > self.capacity:
-            self.remove_oldest()
-
     def popitem(self, remove_pinned: bool = False):
         """Remove and return the `(key, value)` pair least recently used."""
         if not remove_pinned:

vllm/utils/multi_stream_utils.py

@@ -8,10 +8,6 @@ from typing import Any
 import torch
 
 
-class AuxStreamType(Enum):
-    Attention = 1
-
-
 class EventType(Enum):
     Main = 0
     Attention = 1

vllm/v1/core/kv_cache_utils.py

@@ -427,26 +427,6 @@ class FreeKVCacheBlockQueue:
             curr_block = curr_block.next_free_block
 
 
-def need_extra_keys(request: Request) -> bool:
-    """Check whether the blocks allocated to this request need extra hash keys.
-
-    Args:
-        request (Request): The request.
-
-    Returns:
-        bool: Whether blocks allocated to this request need extra hash keys.
-    """
-
-    # Multimodal requests need to include the MM hash.
-    # LoRA requests need to include the LoRA name.
-    # Request with provided cache salt need to include the salt.
-    return (
-        bool(request.mm_features)
-        or (request.lora_request is not None)
-        or (request.cache_salt is not None)
-    )
-
-
 def _gen_mm_extra_hash_keys(
     request: Request, start_token_idx: int, end_token_idx: int, start_mm_idx: int
 ) -> tuple[list[Any], int]:
@@ -1053,20 +1033,6 @@ def _get_kv_cache_groups_uniform_type(
     return [KVCacheGroupSpec(list(spec.kv_cache_specs.keys()), spec)]
 
 
-def is_kv_cache_page_size_uniform(kv_cache_spec: dict[str, KVCacheSpec]) -> bool:
-    """
-    Whether all layers in the given KVCacheSpec have the same page size.
-    Args:
-        kv_cache_spec: The KVCacheSpec of each attention layer in the model
-
-    Returns:
-        True if all layers have the same page size, False otherwise.
-    """
-
-    page_sizes = {layer.page_size_bytes for layer in kv_cache_spec.values()}
-    return len(page_sizes) == 1
-
-
 def unify_kv_cache_spec_page_size(
     kv_cache_spec: dict[str, KVCacheSpec],
 ) -> dict[str, KVCacheSpec]: