Fix & linting

Signed-off-by: Zhanda <zhandazhu@gmail.com>

Author: zhandaz

vllm/envs.py

@@ -1456,6 +1456,7 @@ def get_vllm_port() -> int | None:
     # average at BF16; 100 entries ≈ 0.9 GB.
     "VLLM_POS_EMBED_CACHE_SIZE": lambda: int(
         os.getenv("VLLM_POS_EMBED_CACHE_SIZE", "100")
+    ),
     # Controls whether to use FP8 attention for multimodal encoder (e.g., ViT)
     "VLLM_MM_ENCODER_FP8_ATTN": lambda: bool(
         int(os.getenv("VLLM_MM_ENCODER_FP8_ATTN", "0"))

vllm/model_executor/layers/attention/mm_encoder_attention.py

@@ -53,7 +53,7 @@ def _load_fp8_scales_file(path: str | None) -> dict[str, dict[str, float]]:
         if q is not None and k is not None and v is not None:
             scales[layer_name] = {"q": float(q), "k": float(k), "v": float(v)}
 
-    logger.info(f"Loaded FP8 attention scales from {path} ({len(scales)} layers)")
+    logger.info("Loaded FP8 attention scales from %s (%d layers)", path, len(scales))
     return scales
 
 
@@ -180,12 +180,13 @@ def _init_fp8_attention(self, layer_name: str) -> None:
         self.skip_scale_q = self.fp8_scales["q"] == 1.0
         self.skip_scale_k = self.fp8_scales["k"] == 1.0
         self.skip_scale_v = self.fp8_scales["v"] == 1.0
-        
+
         logger.debug(
-            f"FP8 attention enabled for {layer_name}: "
-            f"q={self.fp8_scales['q']:.4f}, "
-            f"k={self.fp8_scales['k']:.4f}, "
-            f"v={self.fp8_scales['v']:.4f}"
+            "FP8 attention enabled for %s: q=%.4f, k=%.4f, v=%.4f",
+            layer_name,
+            self.fp8_scales["q"],
+            self.fp8_scales["k"],
+            self.fp8_scales["v"],
         )
 
     @classmethod
@@ -292,7 +293,7 @@ def _quantize_to_fp8(
         skip_scale: bool = False,
     ) -> torch.Tensor:
         """Quantize a 3D (S, H, D) tensor to FP8.
-        
+
         Uses QuantFP8 CustomOp when head_dim is aligned to 16; otherwise
         falls back to a stride-aware Triton kernel that pads head_dim to
         a multiple of 16 — no extra copy even for non-contiguous inputs.
@@ -307,9 +308,7 @@ def _quantize_to_fp8(
 
             # QuantFP8 expects 2D input: (total_tokens, num_heads * head_dim)
             tensor_2d = tensor.reshape(orig_shape[0], -1)
-            fp8_tensor, _ = self.fp8_quant.forward_cuda(
-                tensor_2d, scale=scale
-            )
+            fp8_tensor, _ = self.fp8_quant.forward_cuda(tensor_2d, scale=scale)
             return fp8_tensor.reshape(orig_shape)
 
         # Fall back to Triton kernel for padding head_dim to a multiple of 16
@@ -327,9 +326,15 @@ def _forward_flashinfer(
     ) -> torch.Tensor:
         if self.fp8_enabled:
             assert self.fp8_quant is not None and self.fp8_scales is not None
-            query = self._quantize_to_fp8(query, self._fp8_q_scale, skip_scale=self.skip_scale_q)
-            key = self._quantize_to_fp8(key, self._fp8_k_scale, skip_scale=self.skip_scale_k)
-            value = self._quantize_to_fp8(value, self._fp8_v_scale, skip_scale=self.skip_scale_v)
+            query = self._quantize_to_fp8(
+                query, self._fp8_q_scale, skip_scale=self.skip_scale_q
+            )
+            key = self._quantize_to_fp8(
+                key, self._fp8_k_scale, skip_scale=self.skip_scale_k
+            )
+            value = self._quantize_to_fp8(
+                value, self._fp8_v_scale, skip_scale=self.skip_scale_v
+            )
 
         output = vit_flashinfer_wrapper(
             q=query,
@@ -348,7 +353,7 @@ def _forward_flashinfer(
 
         # Un-pad head dimension if it was padded during FP8 quantization
         if self.fp8_enabled and output.shape[-1] != self.head_size:
-            output = output[..., :self.head_size]
+            output = output[..., : self.head_size]
             output = output.contiguous()
 
         return output

vllm/model_executor/layers/quantization/input_quant_fp8.py

@@ -2,6 +2,7 @@
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 
 import os
+
 import torch
 import torch.nn.functional as F
 
@@ -33,7 +34,7 @@ def _quantize_pad_fp8_kernel(
     stride_yh,  # output stride along head dim
     stride_yd,  # output stride along head_dim dim (usually 1)
     num_heads,
-    n_rows,     # total rows = S * H
+    n_rows,  # total rows = S * H
     n_cols,
     n_cols_padded,
     fp8_min,
@@ -56,10 +57,12 @@ def _quantize_pad_fp8_kernel(
     s = offs_m // num_heads
     h = offs_m % num_heads
 
-    x_ptrs = (x_ptr
-              + s[:, None] * stride_xs
-              + h[:, None] * stride_xh
-              + offs_n[None, :] * stride_xd)
+    x_ptrs = (
+        x_ptr
+        + s[:, None] * stride_xs
+        + h[:, None] * stride_xh
+        + offs_n[None, :] * stride_xd
+    )
     x = tl.load(x_ptrs, mask=mask_in, other=0.0).to(tl.float32)
     if SKIP_SCALE:
         x_q = x
@@ -69,10 +72,12 @@ def _quantize_pad_fp8_kernel(
     x_q = tl.where(mask_in, x_q, 0.0)
     x_q = tl.clamp(x_q, fp8_min, fp8_max).to(y_ptr.dtype.element_ty)
 
-    y_ptrs = (y_ptr
-              + s[:, None] * stride_ys
-              + h[:, None] * stride_yh
-              + offs_n[None, :] * stride_yd)
+    y_ptrs = (
+        y_ptr
+        + s[:, None] * stride_ys
+        + h[:, None] * stride_yh
+        + offs_n[None, :] * stride_yd
+    )
     tl.store(y_ptrs, x_q, mask=mask_out)
 
 
@@ -107,24 +112,20 @@ def quantize_fp8_pad_head_dim_triton(
     num_warps: int | None = None,
     block_m: int | None = None,
 ) -> torch.Tensor:
-    """Quantize a 4D (B, S, H, D) or 3D (S, H, D) tensor to FP8 while padding D to a multiple of 16.
+    """Quantize a 3D/4D tensor to FP8, padding head_dim to a multiple of 16.
 
     Reads directly from the input using its 3D strides, so non-contiguous
     views (e.g. Q/K/V slices from an interleaved QKV buffer) are handled
     without an extra copy.  Output is always a fresh contiguous tensor
     with shape (S, H, padded_D).
     """
     if not HAS_TRITON:
-        raise RuntimeError(
-            "Triton is required to quantize with head_dim padding."
-        )
+        raise RuntimeError("Triton is required to quantize with head_dim padding.")
 
     original_shape = tensor.shape
     if tensor.dim() == 4:
         tensor = tensor.view(-1, tensor.shape[-2], tensor.shape[-1])
-    assert tensor.dim() == 3, (
-        f"Expected 3D input (S, H, D), got {tensor.dim()}D"
-    )
+    assert tensor.dim() == 3, f"Expected 3D input (S, H, D), got {tensor.dim()}D"
     S, H, D = tensor.shape
     padded_head_dim = (D + 15) // 16 * 16
     out_dtype = current_platform.fp8_dtype()
@@ -141,15 +142,18 @@ def quantize_fp8_pad_head_dim_triton(
     if block_n is None or num_warps is None or block_m is None:
         block_n, num_warps, block_m = _get_fp8_pad_quant_config(padded_head_dim)
 
-    grid = (triton.cdiv(n_rows, block_m),
-            triton.cdiv(padded_head_dim, block_n))
+    grid = (triton.cdiv(n_rows, block_m), triton.cdiv(padded_head_dim, block_n))
 
     _quantize_pad_fp8_kernel[grid](
         tensor,
         output,
         scale_1d,
-        tensor.stride(0), tensor.stride(1), tensor.stride(2),
-        output.stride(0), output.stride(1), output.stride(2),
+        tensor.stride(0),
+        tensor.stride(1),
+        tensor.stride(2),
+        output.stride(0),
+        output.stride(1),
+        output.stride(2),
         H,
         n_rows,
         D,