[Perf] Trim redundant per-step Python work on the GDN decode and V2 spec-decode paths

vllm/model_executor/layers/mamba/gdn/qwen_gdn_linear_attn.py

@@ -1331,9 +1331,8 @@ def _forward_core(
         a = a[:num_actual_tokens]
 
         # 1. Convolution sequence transformation
-        conv_weights = self.conv1d.weight.view(
-            self.conv1d.weight.size(0), self.conv1d.weight.size(2)
-        )
+        w = self.conv1d.weight
+        conv_weights = w.view(w.size(0), w.size(2))
 
         if spec_sequence_masks is not None:
             if attn_metadata.num_prefills == 0 and attn_metadata.num_decodes == 0:
@@ -1609,9 +1608,8 @@ def _forward_core_decode_aiter(
         ssm_state = self_kv_cache[1]
 
         # 1. Convolution sequence transformation
-        conv_weights = self.conv1d.weight.view(
-            self.conv1d.weight.size(0), self.conv1d.weight.size(2)
-        )
+        w = self.conv1d.weight
+        conv_weights = w.view(w.size(0), w.size(2))
 
         mixed_qkv_non_spec, b, a = (
             gdn_aiter_fused_reshape_causal_conv1d_update_single_token(
@@ -1681,9 +1679,8 @@ def _forward_core_decode_non_spec(
         b = b[:num_actual_tokens]
         a = a[:num_actual_tokens]
 
-        conv_weights = self.conv1d.weight.view(
-            self.conv1d.weight.size(0), self.conv1d.weight.size(2)
-        )
+        w = self.conv1d.weight
+        conv_weights = w.view(w.size(0), w.size(2))
         mixed_qkv_non_spec = causal_conv1d_update(
             mixed_qkv,
             conv_state,
@@ -1731,9 +1728,8 @@ def _forward_core_decode_spec_fused_norm(
             if is_conv_state_dim_first()
             else self.kv_cache[0].transpose(-1, -2)
         )
-        conv_weights = self.conv1d.weight.view(
-            self.conv1d.weight.size(0), self.conv1d.weight.size(2)
-        )
+        w = self.conv1d.weight
+        conv_weights = w.view(w.size(0), w.size(2))
         mixed_qkv = causal_conv1d_update(
             mixed_qkv[:num_actual_tokens],
             conv_state,

vllm/model_executor/layers/mamba/ops/causal_conv1d.py

@@ -10,6 +10,7 @@
 
 from vllm.platforms import current_platform
 from vllm.triton_utils import tl, triton
+from vllm.utils.math_utils import next_power_of_2
 from vllm.v1.attention.backends.utils import NULL_BLOCK_ID, PAD_SLOT_ID
 
 
@@ -577,7 +578,7 @@ def causal_conv1d_fn(
     dim, cu_seqlen = x.shape
     _, width = weight.shape
     state_len = width - 1
-    np2_statelen = triton.next_power_of_2(state_len)
+    np2_statelen = next_power_of_2(state_len)
 
     padded_batch = query_start_loc.size(0) - 1
     stride_x_dim = x.stride(0)
@@ -756,7 +757,7 @@ def grid(META):
         num_stages=2,
         launch_pdl=current_platform.is_arch_support_pdl(),
     )
-    return out.to(original_x_dtype)
+    return out if out.dtype == original_x_dtype else out.to(original_x_dtype)
 
 
 @triton.jit(do_not_specialize_on_alignment=["num_cache_lines"])
@@ -1157,7 +1158,9 @@ def causal_conv1d_update(
         assert activation in ["silu", "swish"]
 
     original_x_dtype = x.dtype
-    x = x.to(conv_state.dtype)
+    conv_state_dtype = conv_state.dtype
+    if original_x_dtype != conv_state_dtype:
+        x = x.to(conv_state_dtype)
     if out is None:
         out = x
     else:
@@ -1221,7 +1224,7 @@ def causal_conv1d_update(
         state_len = width - 1 + (seqlen - 1)  # effective state_len needed
     else:
         state_len = width - 1
-    np2_statelen = triton.next_power_of_2(state_len)
+    np2_statelen = next_power_of_2(state_len)
 
     def grid(META):
         return (
@@ -1276,7 +1279,7 @@ def grid(META):
     )
     if unsqueeze:
         out = out.squeeze(-1)
-    return out.to(original_x_dtype)
+    return out if out.dtype == original_x_dtype else out.to(original_x_dtype)
 
 
 if current_platform.is_cpu():

vllm/v1/worker/gpu/attn_utils.py

@@ -288,6 +288,13 @@ def build_attn_metadata(
     attn_metadata: dict[str, Any] = {}
     num_kv_cache_groups = len(kv_cache_config.kv_cache_groups)
     for i in range(num_kv_cache_groups):
+        groups = attn_groups[i]
+        if not groups:
+            # This model owns no layer in this KV cache group. A drafter builds
+            # its attention groups from active_layer_names only, so
+            # init_attn_backend appends [] for every group it does not own, and
+            # everything below would be constructed and immediately discarded.
+            continue
         block_table = block_tables[i]
         slot_mapping = slot_mappings[i]
         # Per-group causal for hybrid drafters (mixed SWA/full attention).
@@ -325,7 +332,7 @@ def build_attn_metadata(
             **common_attn_metadata_extra_kwargs,
         )
 
-        for attn_group in attn_groups[i]:
+        for attn_group in groups:
             attn_metadata_builder = attn_group.get_metadata_builder(ubatch_idx)
             if for_cudagraph_capture:
                 metadata = attn_metadata_builder.build_for_cudagraph_capture(

vllm/v1/worker/gpu/model_runner.py

@@ -1243,8 +1243,6 @@ def prepare_inputs(
                 count=num_reqs,
             )
             num_bonus_tokens = self.model_state.num_new_sampled_tokens_per_step
-            total_num_draft_tokens = int(num_draft_tokens_per_req.sum())
-            total_num_logits = num_reqs * num_bonus_tokens + total_num_draft_tokens
             num_logits = num_draft_tokens_per_req + num_bonus_tokens
             # combine_sampled_and_draft_tokens places a request's logits rows
             # at [query_end - num_logits, query_end). Fewer query rows than
@@ -1253,6 +1251,10 @@ def prepare_inputs(
             cu_num_logits_np = np.empty(num_reqs + 1, dtype=np.int32)
             cu_num_logits_np[0] = 0
             np.cumsum(num_logits, out=cu_num_logits_np[1:])
+            # The cumsum's last element IS the total, so the separate .sum()
+            # above was a second pass over the same data.
+            total_num_logits = int(cu_num_logits_np[-1])
+            total_num_draft_tokens = total_num_logits - num_reqs * num_bonus_tokens
             cu_num_logits = async_copy_to_gpu(cu_num_logits_np, device=self.device)
 
         adaptive_verification = (

vllm/v1/worker/gpu/spec_decode/rejection_sampler.py

@@ -211,8 +211,16 @@ def _verify_in_chunks(
         for start, end in request_chunks:
             lo = int(cu_num_logits_np[start])
             hi = int(cu_num_logits_np[end])
-            chunk_cu_num_logits_np = cu_num_logits_np[start : end + 1] - lo
-            chunk_cu_num_logits = input_batch.cu_num_logits[start : end + 1] - lo
+            if lo:
+                chunk_cu_num_logits_np = cu_num_logits_np[start : end + 1] - lo
+                chunk_cu_num_logits = input_batch.cu_num_logits[start : end + 1] - lo
+            else:
+                # cu_num_logits always starts at 0, so the single-chunk case
+                # (and the first chunk of a split batch) rebases by nothing.
+                # Skip the GPU sub: a kernel launch and an allocation for a no-op.
+                # NOTE: these are read-only views of the input batch buffers.
+                chunk_cu_num_logits_np = cu_num_logits_np[start : end + 1]
+                chunk_cu_num_logits = input_batch.cu_num_logits[start : end + 1]
             # draft_logits uses persistent request-state indices and stays global.
             processed_logits, sampled, num_sampled = self._verify(
                 logits[lo:hi],

vllm/v1/worker/gpu/spec_decode/speculator.py

@@ -315,7 +315,9 @@ def _build_draft_attn_metadata(
             ],
             query_start_loc_cpu=query_start_loc_cpu,
             max_query_len=max_query_len,
-            seq_lens=self.input_buffers.seq_lens[:num_reqs_padded],
+            # build_attn_metadata re-applies this exact bound as its first
+            # statement, so pre-slicing here only builds a throwaway view.
+            seq_lens=self.input_buffers.seq_lens,
             dcp_local_seq_lens=(
                 None
                 if dcp_local_seq_lens is None