A high-throughput and memory-efficient inference and serving engine for LLMs
Mirrored from https://github.com/tduncan-lum/Nvidia-Optimization-Stack---VLLM.git
Synced Nov 4, 16:15 UTC
[Refactor] Remove unused dead code (#40640)
Signed-off-by: yewentao256 <zhyanwentao@126.com>
Wentao Ye committed
Apr 24, 2026 at 19:28 UTC
a474da28131f61684849b31e29af0eebaaedc383
4 files changed
+4
-22
vllm/model_executor/layers/batch_invariant.py
+3
-5
@@ -39,7 +39,7 @@ def _matmul_launch_metadata(
39
40
41
@triton.jit
42
-def _compute_pid(tile_id, num_pid_in_group, num_pid_m, GROUP_SIZE_M, NUM_SMS):
42
+def _compute_pid(tile_id, num_pid_in_group, num_pid_m, GROUP_SIZE_M):
43
group_id = tile_id // num_pid_in_group
44
first_pid_m = group_id * GROUP_SIZE_M
45
group_size_m = min(num_pid_m - first_pid_m, GROUP_SIZE_M)
@@ -85,9 +85,7 @@ def matmul_kernel_persistent(
85
num_pid_in_group = GROUP_SIZE_M * num_pid_n
86
87
for tile_id in tl.range(start_pid, num_tiles, NUM_SMS, flatten=True):
88
- pid_m, pid_n = _compute_pid(
89
- tile_id, num_pid_in_group, num_pid_m, GROUP_SIZE_M, NUM_SMS
90
- )
88
+ pid_m, pid_n = _compute_pid(tile_id, num_pid_in_group, num_pid_m, GROUP_SIZE_M)
89
start_m = pid_m * BLOCK_SIZE_M
90
start_n = pid_n * BLOCK_SIZE_N
91
offs_am = start_m + tl.arange(0, BLOCK_SIZE_M)
@@ -124,7 +122,7 @@ def matmul_kernel_persistent(
122
123
tile_id_c += NUM_SMS
124
pid_m, pid_n = _compute_pid(
127
- tile_id_c, num_pid_in_group, num_pid_m, GROUP_SIZE_M, NUM_SMS
125
+ tile_id_c, num_pid_in_group, num_pid_m, GROUP_SIZE_M
126
)
127
offs_cm = pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)
128
offs_cn = pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)
vllm/v1/worker/dp_utils.py
-5
@@ -1,8 +1,6 @@
1
# SPDX-License-Identifier: Apache-2.0
2
# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
3
4
-
5
-import numpy as np
4
import torch
5
import torch.distributed as dist
6
@@ -167,7 +165,6 @@ def coordinate_batch_across_dp(
165
parallel_config: ParallelConfig,
166
num_tokens_padded: int | None = None,
167
uniform_decode: bool | None = None,
170
- num_scheduled_tokens_per_request: np.ndarray | None = None,
168
cudagraph_mode: int = 0,
169
) -> tuple[bool, torch.Tensor | None, int]:
170
"""
@@ -182,8 +179,6 @@ def coordinate_batch_across_dp(
179
TP, etc)
180
uniform_decode: Only used if allow_microbatching is True. True if the batch
181
only contains single token decodes
185
- num_scheduled_tokens_per_request: Only used if allow_microbatching is True. The
186
- number of tokens per request.
182
cudagraph_mode: The cudagraph mode for this rank (0=NONE, 1=PIECEWISE, 2=FULL).
183
DP padding is enabled when synced cudagraph mode across ranks is not NONE.
184
vllm/v1/worker/gpu/dp_utils.py
-6
@@ -13,12 +13,6 @@ from vllm.v1.worker.gpu.cudagraph_utils import (
13
)
14
15
16
-def make_num_tokens_across_dp(dp_size: int, num_tokens: int) -> torch.Tensor | None:
17
- if dp_size == 1:
18
- return None
19
- return torch.full((dp_size,), num_tokens, dtype=torch.int32, device="cpu")
20
-
21
-
16
def sync_cudagraph_and_dp_padding(
17
cudagraph_manager: CudaGraphManager | None,
18
desired_batch_desc: BatchExecutionDescriptor,
vllm/v1/worker/gpu_model_runner.py
+1
-6
@@ -3362,7 +3362,6 @@ class GPUModelRunner(
3362
logits: torch.Tensor | None,
3363
hidden_states: torch.Tensor,
3364
num_scheduled_tokens: int,
3365
- spec_decode_metadata: SpecDecodeMetadata | None,
3365
) -> tuple[
3366
dict[str, int],
3367
LogprobsLists | None,
@@ -3630,7 +3629,6 @@ class GPUModelRunner(
3629
allow_microbatching=allow_microbatching,
3630
num_tokens_padded=num_tokens_padded,
3631
uniform_decode=uniform_decode,
3633
- num_scheduled_tokens_per_request=num_scheduled_tokens_np,
3632
cudagraph_mode=cudagraph_mode.value,
3633
)
3634
)
@@ -4308,7 +4306,6 @@ class GPUModelRunner(
4306
logits,
4307
hidden_states,
4308
scheduler_output.total_num_scheduled_tokens,
4311
- spec_decode_metadata,
4309
)
4310
4311
if propose_drafts_after_bookkeeping:
@@ -6540,7 +6537,6 @@ class GPUModelRunner(
6537
6538
def _reshape_kv_cache_tensors(
6539
self,
6543
- kv_cache_config: KVCacheConfig,
6540
kv_cache_raw_tensors: dict[str, torch.Tensor],
6541
kernel_block_sizes: list[int],
6542
) -> dict[str, torch.Tensor]:
@@ -6548,7 +6544,6 @@ class GPUModelRunner(
6544
Reshape the KV cache tensors to the desired shape and dtype.
6545
6546
Args:
6551
- kv_cache_config: The KV cache config
6547
kv_cache_raw_tensors: The KV cache buffer of each layer, with
6548
correct size but uninitialized shape.
6549
kernel_block_sizes: The kernel block sizes for each KV cache group.
@@ -6712,7 +6707,7 @@ class GPUModelRunner(
6707
6708
# Change the memory buffer to the desired shape
6709
kv_caches = self._reshape_kv_cache_tensors(
6715
- kv_cache_config, kv_cache_raw_tensors, kernel_block_sizes
6710
+ kv_cache_raw_tensors, kernel_block_sizes
6711
)
6712
6713
# Set up cross-layer KV cache sharing