add git

Signed-off-by: khluu <khluu000@gmail.com>

benchmarks/kernels/benchmark_moe.py

@@ -226,10 +226,9 @@ def benchmark_config(
                 x, input_gating, topk, renormalize=not use_deep_gemm
             )
 
-            inplace = not disable_inplace()
             if use_deep_gemm:
                 return deep_gemm_experts(
-                    x, w1, w2, topk_weights, topk_ids, inplace=inplace
+                    x, w1, w2, topk_weights, topk_ids, inplace=True
                 )
             return fused_experts(
                 x,
@@ -237,7 +236,7 @@ def benchmark_config(
                 w2,
                 topk_weights,
                 topk_ids,
-                inplace=inplace,
+                inplace=True,
                 quant_config=quant_config,
             )
 

csrc/activation_kernels.cu

@@ -9,111 +9,6 @@
 
 namespace vllm {
 
-struct alignas(32) u32x8_t {
-  uint32_t u0, u1, u2, u3, u4, u5, u6, u7;
-};
-
-__device__ __forceinline__ void ld256(u32x8_t& val, const u32x8_t* ptr) {
-#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 1000
-  asm volatile("ld.global.nc.v8.u32 {%0,%1,%2,%3,%4,%5,%6,%7}, [%8];\n"
-               : "=r"(val.u0), "=r"(val.u1), "=r"(val.u2), "=r"(val.u3),
-                 "=r"(val.u4), "=r"(val.u5), "=r"(val.u6), "=r"(val.u7)
-               : "l"(ptr));
-#else
-  const uint4* uint_ptr = reinterpret_cast<const uint4*>(ptr);
-  uint4 top_half = __ldg(&uint_ptr[0]);
-  uint4 bottom_half = __ldg(&uint_ptr[1]);
-  val.u0 = top_half.x;
-  val.u1 = top_half.y;
-  val.u2 = top_half.z;
-  val.u3 = top_half.w;
-  val.u4 = bottom_half.x;
-  val.u5 = bottom_half.y;
-  val.u6 = bottom_half.z;
-  val.u7 = bottom_half.w;
-#endif
-}
-
-__device__ __forceinline__ void st256(u32x8_t& val, u32x8_t* ptr) {
-#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 1000
-  asm volatile("st.global.v8.u32 [%0], {%1,%2,%3,%4,%5,%6,%7,%8};\n"
-               :
-               : "l"(ptr), "r"(val.u0), "r"(val.u1), "r"(val.u2), "r"(val.u3),
-                 "r"(val.u4), "r"(val.u5), "r"(val.u6), "r"(val.u7)
-               : "memory");
-#else
-  uint4* uint_ptr = reinterpret_cast<uint4*>(ptr);
-  uint_ptr[0] = make_uint4(val.u0, val.u1, val.u2, val.u3);
-  uint_ptr[1] = make_uint4(val.u4, val.u5, val.u6, val.u7);
-#endif
-}
-
-template <bool support_256>
-struct VecTraits;
-
-template <>
-struct VecTraits<true> {
-  static constexpr int ARCH_MAX_VEC_SIZE = 32;
-  using vec_t = u32x8_t;
-};
-
-template <>
-struct VecTraits<false> {
-  static constexpr int ARCH_MAX_VEC_SIZE = 16;
-  using vec_t = int4;
-};
-
-template <typename T>
-struct PackedTraits;
-
-template <>
-struct PackedTraits<c10::BFloat16> {
-  using packed_t = __nv_bfloat162;
-};
-
-template <>
-struct PackedTraits<c10::Half> {
-  using packed_t = __half2;
-};
-
-template <>
-struct PackedTraits<float> {
-  using packed_t = float2;
-};
-
-template <typename packed_t>
-__device__ __forceinline__ float2 cast_to_float2(const packed_t& val) {
-  if constexpr (std::is_same_v<packed_t, __nv_bfloat162>) {
-    return __bfloat1622float2(val);
-  } else if constexpr (std::is_same_v<packed_t, __half2>) {
-    return __half22float2(val);
-  } else if constexpr (std::is_same_v<packed_t, float2>) {
-    return float2(val);
-  }
-}
-
-template <typename packed_t>
-__device__ __forceinline__ packed_t cast_to_packed(const float2& val) {
-  if constexpr (std::is_same_v<packed_t, __nv_bfloat162>) {
-    return __float22bfloat162_rn(val);
-  } else if constexpr (std::is_same_v<packed_t, __half2>) {
-    return __float22half2_rn(val);
-  } else if constexpr (std::is_same_v<packed_t, float2>) {
-    return float2(val);
-  }
-}
-
-template <typename packed_t>
-__device__ __forceinline__ packed_t packed_mul(const packed_t& x,
-                                               const packed_t& y) {
-  if constexpr (std::is_same_v<packed_t, __nv_bfloat162> ||
-                std::is_same_v<packed_t, __half2>) {
-    return __hmul2(x, y);
-  } else if constexpr (std::is_same_v<packed_t, float2>) {
-    return make_float2(x.x * y.x, x.y * y.y);
-  }
-}
-
 template <typename scalar_t, scalar_t (*ACT_FN)(const scalar_t&),
           bool act_first>
 __device__ __forceinline__ scalar_t compute(const scalar_t& x,
@@ -121,69 +16,52 @@ __device__ __forceinline__ scalar_t compute(const scalar_t& x,
   return act_first ? ACT_FN(x) * y : x * ACT_FN(y);
 }
 
-template <typename packed_t, packed_t (*PACKED_ACT_FN)(const packed_t&),
-          bool act_first>
-__device__ __forceinline__ packed_t packed_compute(const packed_t& x,
-                                                   const packed_t& y) {
-  return act_first ? packed_mul(PACKED_ACT_FN(x), y)
-                   : packed_mul(x, PACKED_ACT_FN(y));
-}
-
 // Check if all pointers are 16-byte aligned for int4 vectorized access
-__host__ __device__ __forceinline__ bool is_16byte_aligned(const void* ptr) {
+__device__ __forceinline__ bool is_16byte_aligned(const void* ptr) {
   return (reinterpret_cast<uintptr_t>(ptr) & 15) == 0;
 }
 
-// Check if all pointers are 16-byte aligned for longlong4_32a vectorized access
-__host__ __device__ __forceinline__ bool is_32byte_aligned(const void* ptr) {
-  return (reinterpret_cast<uintptr_t>(ptr) & 31) == 0;
-}
-
 // Activation and gating kernel template.
-template <typename scalar_t, typename packed_t,
-          scalar_t (*ACT_FN)(const scalar_t&),
-          packed_t (*PACKED_ACT_FN)(const packed_t&), bool act_first,
-          bool use_vec, bool use_256b = false>
+template <typename scalar_t, scalar_t (*ACT_FN)(const scalar_t&),
+          bool act_first>
 __global__ void act_and_mul_kernel(
     scalar_t* __restrict__ out,          // [..., d]
     const scalar_t* __restrict__ input,  // [..., 2, d]
     const int d) {
-  const scalar_t* x_ptr = input + blockIdx.x * 2 * d;
+  constexpr int VEC_SIZE = 16 / sizeof(scalar_t);
+  const int64_t token_idx = blockIdx.x;
+  const scalar_t* x_ptr = input + token_idx * 2 * d;
   const scalar_t* y_ptr = x_ptr + d;
-  scalar_t* out_ptr = out + blockIdx.x * d;
+  scalar_t* out_ptr = out + token_idx * d;
 
-  if constexpr (use_vec) {
-    // Fast path: 128-bit/256-bit vectorized loop
-    using vec_t = typename VecTraits<use_256b>::vec_t;
-    constexpr int ARCH_MAX_VEC_SIZE = VecTraits<use_256b>::ARCH_MAX_VEC_SIZE;
-    constexpr int VEC_SIZE = ARCH_MAX_VEC_SIZE / sizeof(packed_t);
+  // Check alignment for 128-bit vectorized access.
+  // All three pointers must be 16-byte aligned for safe int4 operations.
+  const bool aligned = is_16byte_aligned(x_ptr) && is_16byte_aligned(y_ptr) &&
+                       is_16byte_aligned(out_ptr);
 
-    const vec_t* x_vec = reinterpret_cast<const vec_t*>(x_ptr);
-    const vec_t* y_vec = reinterpret_cast<const vec_t*>(y_ptr);
-    vec_t* out_vec = reinterpret_cast<vec_t*>(out_ptr);
-    const int num_vecs = d / 2 / VEC_SIZE;
+  if (aligned && d >= VEC_SIZE) {
+    // Fast path: 128-bit vectorized loop
+    const int4* x_vec = reinterpret_cast<const int4*>(x_ptr);
+    const int4* y_vec = reinterpret_cast<const int4*>(y_ptr);
+    int4* out_vec = reinterpret_cast<int4*>(out_ptr);
+    const int num_vecs = d / VEC_SIZE;
+    const int vec_end = num_vecs * VEC_SIZE;
 
     for (int i = threadIdx.x; i < num_vecs; i += blockDim.x) {
-      vec_t x, y;
-      if constexpr (use_256b) {
-        ld256(x, &x_vec[i]);
-        ld256(y, &y_vec[i]);
-      } else {
-        x = VLLM_LDG(&x_vec[i]);
-        y = VLLM_LDG(&y_vec[i]);
-      }
-      auto* xp = reinterpret_cast<packed_t*>(&x);
-      auto* yp = reinterpret_cast<packed_t*>(&y);
+      int4 x = VLLM_LDG(&x_vec[i]), y = VLLM_LDG(&y_vec[i]), r;
+      auto* xp = reinterpret_cast<scalar_t*>(&x);
+      auto* yp = reinterpret_cast<scalar_t*>(&y);
+      auto* rp = reinterpret_cast<scalar_t*>(&r);
 #pragma unroll
       for (int j = 0; j < VEC_SIZE; j++) {
-        xp[j] =
-            packed_compute<packed_t, PACKED_ACT_FN, act_first>(xp[j], yp[j]);
-      }
-      if constexpr (use_256b) {
-        st256(x, &out_vec[i]);
-      } else {
-        out_vec[i] = x;
+        rp[j] = compute<scalar_t, ACT_FN, act_first>(xp[j], yp[j]);
       }
+      out_vec[i] = r;
+    }
+    // Scalar cleanup for remaining elements
+    for (int i = vec_end + threadIdx.x; i < d; i += blockDim.x) {
+      out_ptr[i] = compute<scalar_t, ACT_FN, act_first>(VLLM_LDG(&x_ptr[i]),
+                                                        VLLM_LDG(&y_ptr[i]));
     }
   } else {
     // Scalar fallback for unaligned data or small d
@@ -201,15 +79,6 @@ __device__ __forceinline__ T silu_kernel(const T& x) {
   return (T)(((float)x) / (1.0f + expf((float)-x)));
 }
 
-template <typename packed_t>
-__device__ __forceinline__ packed_t packed_silu_kernel(const packed_t& val) {
-  // x * sigmoid(x)
-  float2 fval = cast_to_float2(val);
-  fval.x = fval.x / (1.0f + expf(-fval.x));
-  fval.y = fval.y / (1.0f + expf(-fval.y));
-  return cast_to_packed<packed_t>(fval);
-}
-
 template <typename T>
 __device__ __forceinline__ T gelu_kernel(const T& x) {
   // Equivalent to PyTorch GELU with 'none' approximation.
@@ -220,18 +89,6 @@ __device__ __forceinline__ T gelu_kernel(const T& x) {
   return (T)(f * 0.5f * (1.0f + ::erf(f * ALPHA)));
 }
 
-template <typename packed_t>
-__device__ __forceinline__ packed_t packed_gelu_kernel(const packed_t& val) {
-  // Equivalent to PyTorch GELU with 'none' approximation.
-  // Refer to:
-  // https://github.com/pytorch/pytorch/blob/8ac9b20d4b090c213799e81acf48a55ea8d437d6/aten/src/ATen/native/cuda/ActivationGeluKernel.cu#L36-L38
-  constexpr float ALPHA = M_SQRT1_2;
-  float2 fval = cast_to_float2(val);
-  fval.x = fval.x * 0.5f * (1.0f + ::erf(fval.x * ALPHA));
-  fval.y = fval.y * 0.5f * (1.0f + ::erf(fval.y * ALPHA));
-  return cast_to_packed<packed_t>(fval);
-}
-
 template <typename T>
 __device__ __forceinline__ T gelu_tanh_kernel(const T& x) {
   // Equivalent to PyTorch GELU with 'tanh' approximation.
@@ -245,83 +102,32 @@ __device__ __forceinline__ T gelu_tanh_kernel(const T& x) {
   return (T)(0.5f * f * (1.0f + ::tanhf(inner)));
 }
 
-template <typename packed_t>
-__device__ __forceinline__ packed_t
-packed_gelu_tanh_kernel(const packed_t& val) {
-  // Equivalent to PyTorch GELU with 'tanh' approximation.
-  // Refer to:
-  // https://github.com/pytorch/pytorch/blob/8ac9b20d4b090c213799e81acf48a55ea8d437d6/aten/src/ATen/native/cuda/ActivationGeluKernel.cu#L25-L30
-  float2 fval = cast_to_float2(val);
-  constexpr float BETA = M_SQRT2 * M_2_SQRTPI * 0.5f;
-  constexpr float KAPPA = 0.044715;
-
-  float x_cube = fval.x * fval.x * fval.x;
-  float inner = BETA * (fval.x + KAPPA * x_cube);
-  fval.x = 0.5f * fval.x * (1.0f + ::tanhf(inner));
-
-  x_cube = fval.y * fval.y * fval.y;
-  inner = BETA * (fval.y + KAPPA * x_cube);
-  fval.y = 0.5f * fval.y * (1.0f + ::tanhf(inner));
-  return cast_to_packed<packed_t>(fval);
-}
-
 }  // namespace vllm
 
 // Launch activation and gating kernel.
 // Use ACT_FIRST (bool) indicating whether to apply the activation function
 // first.
-#define LAUNCH_ACTIVATION_GATE_KERNEL(KERNEL, PACKED_KERNEL, ACT_FIRST)     \
-  auto dtype = input.scalar_type();                                         \
-  int d = input.size(-1) / 2;                                               \
-  int64_t num_tokens = input.numel() / input.size(-1);                      \
-  if (num_tokens == 0) {                                                    \
-    return;                                                                 \
-  }                                                                         \
-  dim3 grid(num_tokens);                                                    \
-  int cc_major = at::cuda::getCurrentDeviceProperties()->major;             \
-  int support_vec = (cc_major >= 10 && num_tokens > 128) ? 32 : 16;         \
-  int vec_size = support_vec / at::elementSize(dtype);                      \
-  const bool use_vec = (d % vec_size == 0);                                 \
-  const at::cuda::OptionalCUDAGuard device_guard(device_of(input));         \
-  const cudaStream_t stream = at::cuda::getCurrentCUDAStream();             \
-  if (use_vec) {                                                            \
-    dim3 block(std::min(d / vec_size, 1024));                               \
-    if (cc_major >= 10 && num_tokens > 128) {                               \
-      VLLM_DISPATCH_FLOATING_TYPES(dtype, "act_and_mul_kernel", [&] {       \
-        vllm::act_and_mul_kernel<                                           \
-            scalar_t, typename vllm::PackedTraits<scalar_t>::packed_t,      \
-            KERNEL<scalar_t>,                                               \
-            PACKED_KERNEL<typename vllm::PackedTraits<scalar_t>::packed_t>, \
-            ACT_FIRST, true, true><<<grid, block, 0, stream>>>(             \
-            out.data_ptr<scalar_t>(), input.data_ptr<scalar_t>(), d);       \
-      });                                                                   \
-    } else {                                                                \
-      VLLM_DISPATCH_FLOATING_TYPES(dtype, "act_and_mul_kernel", [&] {       \
-        vllm::act_and_mul_kernel<                                           \
-            scalar_t, typename vllm::PackedTraits<scalar_t>::packed_t,      \
-            KERNEL<scalar_t>,                                               \
-            PACKED_KERNEL<typename vllm::PackedTraits<scalar_t>::packed_t>, \
-            ACT_FIRST, true, false><<<grid, block, 0, stream>>>(            \
-            out.data_ptr<scalar_t>(), input.data_ptr<scalar_t>(), d);       \
-      });                                                                   \
-    }                                                                       \
-  } else {                                                                  \
-    dim3 block(std::min(d, 1024));                                          \
-    VLLM_DISPATCH_FLOATING_TYPES(dtype, "act_and_mul_kernel", [&] {         \
-      vllm::act_and_mul_kernel<                                             \
-          scalar_t, typename vllm::PackedTraits<scalar_t>::packed_t,        \
-          KERNEL<scalar_t>,                                                 \
-          PACKED_KERNEL<typename vllm::PackedTraits<scalar_t>::packed_t>,   \
-          ACT_FIRST, false><<<grid, block, 0, stream>>>(                    \
-          out.data_ptr<scalar_t>(), input.data_ptr<scalar_t>(), d);         \
-    });                                                                     \
-  }
+#define LAUNCH_ACTIVATION_GATE_KERNEL(KERNEL, ACT_FIRST)                 \
+  int d = input.size(-1) / 2;                                            \
+  int64_t num_tokens = input.numel() / input.size(-1);                   \
+  dim3 grid(num_tokens);                                                 \
+  dim3 block(std::min(d, 1024));                                         \
+  if (num_tokens == 0) {                                                 \
+    return;                                                              \
+  }                                                                      \
+  const at::cuda::OptionalCUDAGuard device_guard(device_of(input));      \
+  const cudaStream_t stream = at::cuda::getCurrentCUDAStream();          \
+  VLLM_DISPATCH_FLOATING_TYPES(                                          \
+      input.scalar_type(), "act_and_mul_kernel", [&] {                   \
+        vllm::act_and_mul_kernel<scalar_t, KERNEL<scalar_t>, ACT_FIRST>  \
+            <<<grid, block, 0, stream>>>(out.data_ptr<scalar_t>(),       \
+                                         input.data_ptr<scalar_t>(), d); \
+      });
 
 void silu_and_mul(torch::Tensor& out,    // [..., d]
                   torch::Tensor& input)  // [..., 2 * d]
 {
-  LAUNCH_ACTIVATION_GATE_KERNEL(vllm::silu_kernel, vllm::packed_silu_kernel,
-                                true);
+  LAUNCH_ACTIVATION_GATE_KERNEL(vllm::silu_kernel, true);
 }
 
 void mul_and_silu(torch::Tensor& out,    // [..., d]
@@ -329,22 +135,19 @@ void mul_and_silu(torch::Tensor& out,    // [..., d]
 {
   // The difference between mul_and_silu and silu_and_mul is that mul_and_silu
   // applies the silu to the latter half of the input.
-  LAUNCH_ACTIVATION_GATE_KERNEL(vllm::silu_kernel, vllm::packed_silu_kernel,
-                                false);
+  LAUNCH_ACTIVATION_GATE_KERNEL(vllm::silu_kernel, false);
 }
 
 void gelu_and_mul(torch::Tensor& out,    // [..., d]
                   torch::Tensor& input)  // [..., 2 * d]
 {
-  LAUNCH_ACTIVATION_GATE_KERNEL(vllm::gelu_kernel, vllm::packed_gelu_kernel,
-                                true);
+  LAUNCH_ACTIVATION_GATE_KERNEL(vllm::gelu_kernel, true);
 }
 
 void gelu_tanh_and_mul(torch::Tensor& out,    // [..., d]
                        torch::Tensor& input)  // [..., 2 * d]
 {
-  LAUNCH_ACTIVATION_GATE_KERNEL(vllm::gelu_tanh_kernel,
-                                vllm::packed_gelu_tanh_kernel, true);
+  LAUNCH_ACTIVATION_GATE_KERNEL(vllm::gelu_tanh_kernel, true);
 }
 
 namespace vllm {
@@ -355,57 +158,42 @@ __device__ __forceinline__ T fatrelu_kernel(const T& x, const float threshold) {
   return (T)(f > threshold ? f : 0.0f);
 }
 
-template <typename packed_t>
-__device__ __forceinline__ packed_t
-packed_fatrelu_kernel(const packed_t& val, const float threshold) {
-  float2 fval = cast_to_float2(val);
-  fval.x = fval.x > threshold ? fval.x : 0.0f;
-  fval.y = fval.y > threshold ? fval.y : 0.0f;
-  return cast_to_packed<packed_t>(fval);
-}
-
-template <typename scalar_t, typename packed_t,
-          scalar_t (*ACT_FN)(const scalar_t&, const float),
-          packed_t (*PACKED_ACT_FN)(const packed_t&, const float), bool use_vec,
-          bool use_256b = false>
+template <typename scalar_t, scalar_t (*ACT_FN)(const scalar_t&, const float)>
 __global__ void act_and_mul_kernel_with_param(
     scalar_t* __restrict__ out, const scalar_t* __restrict__ input, const int d,
     const float param) {
-  const scalar_t* x_ptr = input + blockIdx.x * 2 * d;
+  constexpr int VEC_SIZE = 16 / sizeof(scalar_t);
+  const int64_t token_idx = blockIdx.x;
+  const scalar_t* x_ptr = input + token_idx * 2 * d;
   const scalar_t* y_ptr = x_ptr + d;
-  scalar_t* out_ptr = out + blockIdx.x * d;
+  scalar_t* out_ptr = out + token_idx * d;
 
-  if constexpr (use_vec) {
-    // Fast path: 128-bit/256-bit vectorized loop
-    using vec_t = typename VecTraits<use_256b>::vec_t;
-    constexpr int ARCH_MAX_VEC_SIZE = VecTraits<use_256b>::ARCH_MAX_VEC_SIZE;
-    constexpr int VEC_SIZE = ARCH_MAX_VEC_SIZE / sizeof(packed_t);
+  // Check alignment for 128-bit vectorized access
+  const bool aligned = is_16byte_aligned(x_ptr) && is_16byte_aligned(y_ptr) &&
+                       is_16byte_aligned(out_ptr);
 
-    const vec_t* x_vec = reinterpret_cast<const vec_t*>(x_ptr);
-    const vec_t* y_vec = reinterpret_cast<const vec_t*>(y_ptr);
-    vec_t* out_vec = reinterpret_cast<vec_t*>(out_ptr);
-    const int num_vecs = d / 2 / VEC_SIZE;
+  if (aligned && d >= VEC_SIZE) {
+    // Fast path: 128-bit vectorized loop
+    const int4* x_vec = reinterpret_cast<const int4*>(x_ptr);
+    const int4* y_vec = reinterpret_cast<const int4*>(y_ptr);
+    int4* out_vec = reinterpret_cast<int4*>(out_ptr);
+    const int num_vecs = d / VEC_SIZE;
+    const int vec_end = num_vecs * VEC_SIZE;
 
     for (int i = threadIdx.x; i < num_vecs; i += blockDim.x) {
-      vec_t x, y;
-      if constexpr (use_256b) {
-        ld256(x, &x_vec[i]);
-        ld256(y, &y_vec[i]);
-      } else {
-        x = VLLM_LDG(&x_vec[i]);
-        y = VLLM_LDG(&y_vec[i]);
-      }
-      auto* xp = reinterpret_cast<packed_t*>(&x);
-      auto* yp = reinterpret_cast<packed_t*>(&y);
+      int4 x = VLLM_LDG(&x_vec[i]), y = VLLM_LDG(&y_vec[i]), r;
+      auto* xp = reinterpret_cast<scalar_t*>(&x);
+      auto* yp = reinterpret_cast<scalar_t*>(&y);
+      auto* rp = reinterpret_cast<scalar_t*>(&r);
 #pragma unroll
       for (int j = 0; j < VEC_SIZE; j++) {
-        xp[j] = packed_mul(PACKED_ACT_FN(xp[j], param), yp[j]);
-      }
-      if constexpr (use_256b) {
-        st256(x, &out_vec[i]);
-      } else {
-        out_vec[i] = x;
+        rp[j] = ACT_FN(xp[j], param) * yp[j];
       }
+      out_vec[i] = r;
+    }
+    // Scalar cleanup for remaining elements
+    for (int i = vec_end + threadIdx.x; i < d; i += blockDim.x) {
+      out_ptr[i] = ACT_FN(VLLM_LDG(&x_ptr[i]), param) * VLLM_LDG(&y_ptr[i]);
     }
   } else {
     // Scalar fallback for unaligned data or small d
@@ -488,58 +276,20 @@ __global__ void swigluoai_and_mul_kernel(
 
 }  // namespace vllm
 
-#define LAUNCH_ACTIVATION_GATE_KERNEL_WITH_PARAM(KERNEL, PACKED_KERNEL, PARAM) \
-  auto dtype = input.scalar_type();                                            \
-  int d = input.size(-1) / 2;                                                  \
-  int64_t num_tokens = input.numel() / input.size(-1);                         \
-  if (num_tokens == 0) {                                                       \
-    return;                                                                    \
-  }                                                                            \
-  dim3 grid(num_tokens);                                                       \
-  int cc_major = at::cuda::getCurrentDeviceProperties()->major;                \
-  int support_vec = (cc_major >= 10 && num_tokens > 128) ? 32 : 16;            \
-  int vec_size = support_vec / at::elementSize(dtype);                         \
-  const bool use_vec = (d % vec_size == 0);                                    \
-  const at::cuda::OptionalCUDAGuard device_guard(device_of(input));            \
-  const cudaStream_t stream = at::cuda::getCurrentCUDAStream();                \
-  if (use_vec) {                                                               \
-    dim3 block(std::min(d / vec_size, 1024));                                  \
-    if (cc_major >= 10 && num_tokens > 128) {                                  \
-      VLLM_DISPATCH_FLOATING_TYPES(                                            \
-          dtype, "act_and_mul_kernel_with_param", [&] {                        \
-            vllm::act_and_mul_kernel_with_param<                               \
-                scalar_t, typename vllm::PackedTraits<scalar_t>::packed_t,     \
-                KERNEL<scalar_t>,                                              \
-                PACKED_KERNEL<                                                 \
-                    typename vllm::PackedTraits<scalar_t>::packed_t>,          \
-                true, true><<<grid, block, 0, stream>>>(                       \
-                out.data_ptr<scalar_t>(), input.data_ptr<scalar_t>(), d,       \
-                PARAM);                                                        \
-          });                                                                  \
-    } else {                                                                   \
-      VLLM_DISPATCH_FLOATING_TYPES(                                            \
-          dtype, "act_and_mul_kernel_with_param", [&] {                        \
-            vllm::act_and_mul_kernel_with_param<                               \
-                scalar_t, typename vllm::PackedTraits<scalar_t>::packed_t,     \
-                KERNEL<scalar_t>,                                              \
-                PACKED_KERNEL<                                                 \
-                    typename vllm::PackedTraits<scalar_t>::packed_t>,          \
-                true, false><<<grid, block, 0, stream>>>(                      \
-                out.data_ptr<scalar_t>(), input.data_ptr<scalar_t>(), d,       \
-                PARAM);                                                        \
-          });                                                                  \
-    }                                                                          \
-  } else {                                                                     \
-    dim3 block(std::min(d, 1024));                                             \
-    VLLM_DISPATCH_FLOATING_TYPES(dtype, "act_and_mul_kernel_with_param", [&] { \
-      vllm::act_and_mul_kernel_with_param<                                     \
-          scalar_t, typename vllm::PackedTraits<scalar_t>::packed_t,           \
-          KERNEL<scalar_t>,                                                    \
-          PACKED_KERNEL<typename vllm::PackedTraits<scalar_t>::packed_t>,      \
-          false><<<grid, block, 0, stream>>>(                                  \
-          out.data_ptr<scalar_t>(), input.data_ptr<scalar_t>(), d, PARAM);     \
-    });                                                                        \
-  }
+#define LAUNCH_ACTIVATION_GATE_KERNEL_WITH_PARAM(KERNEL, PARAM)         \
+  int d = input.size(-1) / 2;                                           \
+  int64_t num_tokens = input.numel() / input.size(-1);                  \
+  dim3 grid(num_tokens);                                                \
+  dim3 block(std::min(d, 1024));                                        \
+  const at::cuda::OptionalCUDAGuard device_guard(device_of(input));     \
+  const cudaStream_t stream = at::cuda::getCurrentCUDAStream();         \
+  VLLM_DISPATCH_FLOATING_TYPES(                                         \
+      input.scalar_type(), "act_and_mul_kernel_with_param", [&] {       \
+        vllm::act_and_mul_kernel_with_param<scalar_t, KERNEL<scalar_t>> \
+            <<<grid, block, 0, stream>>>(out.data_ptr<scalar_t>(),      \
+                                         input.data_ptr<scalar_t>(), d, \
+                                         PARAM);                        \
+      });
 
 #define LAUNCH_SIGLUOAI_AND_MUL(KERNEL, ALPHA, LIMIT)                          \
   int d = input.size(-1) / 2;                                                  \
@@ -559,8 +309,7 @@ __global__ void swigluoai_and_mul_kernel(
 void fatrelu_and_mul(torch::Tensor& out,    // [..., d],
                      torch::Tensor& input,  // [..., 2 * d]
                      double threshold) {
-  LAUNCH_ACTIVATION_GATE_KERNEL_WITH_PARAM(
-      vllm::fatrelu_kernel, vllm::packed_fatrelu_kernel, threshold);
+  LAUNCH_ACTIVATION_GATE_KERNEL_WITH_PARAM(vllm::fatrelu_kernel, threshold);
 }
 void swigluoai_and_mul(torch::Tensor& out,    // [..., d]
                        torch::Tensor& input,  // [..., 2 * d]
@@ -570,41 +319,39 @@ void swigluoai_and_mul(torch::Tensor& out,    // [..., d]
 namespace vllm {
 
 // Element-wise activation kernel template.
-template <typename scalar_t, scalar_t (*ACT_FN)(const scalar_t&), bool use_vec,
-          bool use_256b = false>
+template <typename scalar_t, scalar_t (*ACT_FN)(const scalar_t&)>
 __global__ void activation_kernel(
     scalar_t* __restrict__ out,          // [..., d]
     const scalar_t* __restrict__ input,  // [..., d]
     const int d) {
-  const scalar_t* in_ptr = input + blockIdx.x * d;
-  scalar_t* out_ptr = out + blockIdx.x * d;
+  constexpr int VEC_SIZE = 16 / sizeof(scalar_t);
+  const int64_t token_idx = blockIdx.x;
+  const scalar_t* in_ptr = input + token_idx * d;
+  scalar_t* out_ptr = out + token_idx * d;
 
-  if constexpr (use_vec) {
-    // Fast path: 128-bit/256-bit vectorized loop
-    using vec_t = typename VecTraits<use_256b>::vec_t;
-    constexpr int ARCH_MAX_VEC_SIZE = VecTraits<use_256b>::ARCH_MAX_VEC_SIZE;
-    constexpr int VEC_SIZE = ARCH_MAX_VEC_SIZE / sizeof(scalar_t);
-    const vec_t* in_vec = reinterpret_cast<const vec_t*>(in_ptr);
-    vec_t* out_vec = reinterpret_cast<vec_t*>(out_ptr);
+  // Check alignment for 128-bit vectorized access
+  const bool aligned = is_16byte_aligned(in_ptr) && is_16byte_aligned(out_ptr);
+
+  if (aligned && d >= VEC_SIZE) {
+    // Fast path: 128-bit vectorized loop
+    const int4* in_vec = reinterpret_cast<const int4*>(in_ptr);
+    int4* out_vec = reinterpret_cast<int4*>(out_ptr);
     const int num_vecs = d / VEC_SIZE;
+    const int vec_end = num_vecs * VEC_SIZE;
 
     for (int i = threadIdx.x; i < num_vecs; i += blockDim.x) {
-      vec_t v;
-      if constexpr (use_256b) {
-        ld256(v, &in_vec[i]);
-      } else {
-        v = VLLM_LDG(&in_vec[i]);
-      }
+      int4 v = VLLM_LDG(&in_vec[i]), r;
       auto* vp = reinterpret_cast<scalar_t*>(&v);
+      auto* rp = reinterpret_cast<scalar_t*>(&r);
 #pragma unroll
       for (int j = 0; j < VEC_SIZE; j++) {
-        vp[j] = ACT_FN(vp[j]);
-      }
-      if constexpr (use_256b) {
-        st256(v, &out_vec[i]);
-      } else {
-        out_vec[i] = v;
+        rp[j] = ACT_FN(vp[j]);
       }
+      out_vec[i] = r;
+    }
+    // Scalar cleanup for remaining elements
+    for (int i = vec_end + threadIdx.x; i < d; i += blockDim.x) {
+      out_ptr[i] = ACT_FN(VLLM_LDG(&in_ptr[i]));
     }
   } else {
     // Scalar fallback for unaligned data or small d
@@ -618,43 +365,18 @@ __global__ void activation_kernel(
 }  // namespace vllm
 
 // Launch element-wise activation kernel.
-#define LAUNCH_ACTIVATION_KERNEL(KERNEL)                                 \
-  auto dtype = input.scalar_type();                                      \
-  int d = input.size(-1);                                                \
-  int64_t num_tokens = input.numel() / input.size(-1);                   \
-  if (num_tokens == 0) {                                                 \
-    return;                                                              \
-  }                                                                      \
-  dim3 grid(num_tokens);                                                 \
-  int cc_major = at::cuda::getCurrentDeviceProperties()->major;          \
-  int support_vec = (cc_major >= 10 && num_tokens > 128) ? 32 : 16;      \
-  int vec_size = support_vec / at::elementSize(dtype);                   \
-  const bool use_vec = (d % vec_size == 0);                              \
-  const at::cuda::OptionalCUDAGuard device_guard(device_of(input));      \
-  const cudaStream_t stream = at::cuda::getCurrentCUDAStream();          \
-  if (use_vec) {                                                         \
-    dim3 block(std::min(d / vec_size, 1024));                            \
-    if (cc_major >= 10 && num_tokens > 128) {                            \
-      VLLM_DISPATCH_FLOATING_TYPES(dtype, "activation_kernel", [&] {     \
-        vllm::activation_kernel<scalar_t, KERNEL<scalar_t>, true, true>  \
-            <<<grid, block, 0, stream>>>(out.data_ptr<scalar_t>(),       \
-                                         input.data_ptr<scalar_t>(), d); \
-      });                                                                \
-    } else {                                                             \
-      VLLM_DISPATCH_FLOATING_TYPES(dtype, "activation_kernel", [&] {     \
-        vllm::activation_kernel<scalar_t, KERNEL<scalar_t>, true, false> \
-            <<<grid, block, 0, stream>>>(out.data_ptr<scalar_t>(),       \
-                                         input.data_ptr<scalar_t>(), d); \
-      });                                                                \
-    }                                                                    \
-  } else {                                                               \
-    dim3 block(std::min(d, 1024));                                       \
-    VLLM_DISPATCH_FLOATING_TYPES(dtype, "activation_kernel", [&] {       \
-      vllm::activation_kernel<scalar_t, KERNEL<scalar_t>, false>         \
-          <<<grid, block, 0, stream>>>(out.data_ptr<scalar_t>(),         \
-                                       input.data_ptr<scalar_t>(), d);   \
-    });                                                                  \
-  }
+#define LAUNCH_ACTIVATION_KERNEL(KERNEL)                                       \
+  int d = input.size(-1);                                                      \
+  int64_t num_tokens = input.numel() / d;                                      \
+  dim3 grid(num_tokens);                                                       \
+  dim3 block(std::min(d, 1024));                                               \
+  const at::cuda::OptionalCUDAGuard device_guard(device_of(input));            \
+  const cudaStream_t stream = at::cuda::getCurrentCUDAStream();                \
+  VLLM_DISPATCH_FLOATING_TYPES(input.scalar_type(), "activation_kernel", [&] { \
+    vllm::activation_kernel<scalar_t, KERNEL<scalar_t>>                        \
+        <<<grid, block, 0, stream>>>(out.data_ptr<scalar_t>(),                 \
+                                     input.data_ptr<scalar_t>(), d);           \
+  });
 
 namespace vllm {
 

csrc/cache_kernels.cu

@@ -1234,13 +1234,8 @@ void cp_gather_and_upconvert_fp8_kv_cache(
               "src_cache and seq_lens must be on the same device");
   TORCH_CHECK(src_cache.device() == workspace_starts.device(),
               "src_cache and workspace_starts must be on the same device");
-  auto dtype = src_cache.scalar_type();
-  TORCH_CHECK(
-      dtype == at::ScalarType::Byte ||               // uint8
-          dtype == at::ScalarType::Float8_e4m3fn ||  // fp8 e4m3
-          dtype == at::ScalarType::Float8_e5m2,      // fp8 e5m2
-      "src_cache must be uint8, float8_e4m3fn, or float8_e5m2, but got ",
-      src_cache.dtype());
+
+  TORCH_CHECK(src_cache.dtype() == torch::kUInt8, "src_cache must be uint8");
   TORCH_CHECK(dst.dtype() == torch::kBFloat16, "dst must be bfloat16");
   TORCH_CHECK(head_dim == 576, "head_dim must be 576 for MLA");
 
@@ -1249,21 +1244,14 @@ void cp_gather_and_upconvert_fp8_kv_cache(
   int64_t cache_entry_stride = src_cache.stride(1);
   int64_t dst_entry_stride = dst.stride(0);
 
-  const uint8_t* src_ptr = nullptr;
-  if (dtype == at::ScalarType::Byte) {
-    src_ptr = src_cache.data_ptr<uint8_t>();
-  } else {
-    // float8_e4m3fn or float8_e5m2
-    src_ptr = reinterpret_cast<const uint8_t*>(src_cache.data_ptr());
-  }
-
   // Decide on the number of splits based on the batch size
   int num_splits = batch_size > 128 ? 2 : batch_size > 64 ? 4 : 16;
   dim3 grid(batch_size, num_splits);
   dim3 block(576);
 
   vllm::cp_gather_and_upconvert_fp8_kv_cache<<<grid, block, 0, stream>>>(
-      src_ptr, reinterpret_cast<__nv_bfloat16*>(dst.data_ptr()),
+      src_cache.data_ptr<uint8_t>(),
+      reinterpret_cast<__nv_bfloat16*>(dst.data_ptr()),
       block_table.data_ptr<int32_t>(), seq_lens.data_ptr<int32_t>(),
       workspace_starts.data_ptr<int32_t>(), block_size, head_dim,
       block_table_stride, cache_block_stride, cache_entry_stride,

csrc/cpu/cpu_attn_impl.hpp

@@ -821,7 +821,7 @@ struct VecTypeTrait<c10::BFloat16> {
   using vec_t = vec_op::BF16Vec16;
 };
 
-#if !defined(__powerpc__)
+#if !defined(__powerpc__) && !defined(__s390x__)
 template <>
 struct VecTypeTrait<c10::Half> {
   using vec_t = vec_op::FP16Vec16;

csrc/cpu/cpu_types_vxe.hpp

@@ -16,12 +16,10 @@ namespace vec_op {
 #define vec_sr(a, b) ((a) >> (b))  // Vector Shift Right Algebraic
 #define vec_sl(a, b) ((a) << (b))  // Vector Shift Left
 
-// NOTE: FP16 (Half) is supported on s390x via custom bit-manipulation
-// conversion. PyTorch itself lacks native s390x FP16 support.
-#define VLLM_DISPATCH_CASE_FLOATING_TYPES(...)            \
-  AT_DISPATCH_CASE(at::ScalarType::Float, __VA_ARGS__)    \
-  AT_DISPATCH_CASE(at::ScalarType::BFloat16, __VA_ARGS__) \
-  AT_DISPATCH_CASE(at::ScalarType::Half, __VA_ARGS__)
+// FIXME: FP16 is not fully supported in Torch-CPU
+#define VLLM_DISPATCH_CASE_FLOATING_TYPES(...)         \
+  AT_DISPATCH_CASE(at::ScalarType::Float, __VA_ARGS__) \
+  AT_DISPATCH_CASE(at::ScalarType::BFloat16, __VA_ARGS__)
 
 #define VLLM_DISPATCH_FLOATING_TYPES(TYPE, NAME, ...) \
   AT_DISPATCH_SWITCH(TYPE, NAME, VLLM_DISPATCH_CASE_FLOATING_TYPES(__VA_ARGS__))
@@ -88,39 +86,6 @@ struct BF16Vec8 : public Vec<BF16Vec8> {
   }
 };
 
-struct FP16Vec8 : public Vec<FP16Vec8> {
-  constexpr static int VEC_ELEM_NUM = 8;
-
-  __vector signed short reg;
-
-  explicit FP16Vec8(const void* ptr) : reg(*(__vector signed short*)ptr) {}
-  explicit FP16Vec8(const FP32Vec8&);
-
-  void save(void* ptr) const {
-    *reinterpret_cast<__vector signed short*>(ptr) = reg;
-  }
-};
-
-struct FP16Vec16 : public Vec<FP16Vec16> {
-  constexpr static int VEC_ELEM_NUM = 16;
-
-  ss16x8x2_t reg;
-
-  explicit FP16Vec16(const void* ptr) {
-    // Load 256 bits (16 FP16 values) in two parts
-    reg.val[0] = (__vector signed short)vec_xl(0, (signed short*)ptr);
-    reg.val[1] = (__vector signed short)vec_xl(16, (signed short*)ptr);
-  }
-
-  explicit FP16Vec16(const FP32Vec16&);
-
-  void save(void* ptr) const {
-    // Save 256 bits in two parts
-    vec_xst(reg.val[0], 0, (signed short*)ptr);
-    vec_xst(reg.val[1], 16, (signed short*)ptr);
-  }
-};
-
 struct BF16Vec16 : public Vec<BF16Vec16> {
   constexpr static int VEC_ELEM_NUM = 16;
 
@@ -143,92 +108,6 @@ struct BF16Vec16 : public Vec<BF16Vec16> {
 
 const static __vector signed short zero = vec_splats((signed short)0);
 
-FORCE_INLINE __vector float fp16_to_fp32_bits(__vector unsigned int x) {
-  const __vector unsigned int mask_sign = {0x8000, 0x8000, 0x8000, 0x8000};
-  const __vector unsigned int mask_exp = {0x7C00, 0x7C00, 0x7C00, 0x7C00};
-  const __vector unsigned int mask_mant = {0x03FF, 0x03FF, 0x03FF, 0x03FF};
-  const __vector unsigned int bias_adj = {112, 112, 112, 112};
-  const __vector unsigned int exp_max_fp16 = {0x1F, 0x1F, 0x1F,
-                                              0x1F};  // FP16 NaN/Inf exponent
-  const __vector unsigned int exp_max_fp32 = {0xFF, 0xFF, 0xFF,
-                                              0xFF};  // FP32 NaN/Inf exponent
-
-  __vector unsigned int s = (x & mask_sign) << 16;
-  __vector unsigned int e = (x & mask_exp) >> 10;
-  __vector unsigned int m = (x & mask_mant) << 13;
-
-  // Check for NaN/Inf: exponent = 0x1F in FP16
-  __vector __bool int is_nan_inf = vec_cmpeq(e, exp_max_fp16);
-
-  // Normal: adjust bias; NaN/Inf: set to 0xFF
-  __vector unsigned int e_normal = e + bias_adj;
-  e = vec_sel(e_normal, exp_max_fp32, is_nan_inf);
-
-  return (__vector float)(s | (e << 23) | m);
-}
-
-FORCE_INLINE __vector unsigned int fp32_to_fp16_bits(__vector float f_in) {
-  __vector unsigned int in = (__vector unsigned int)f_in;
-
-  const __vector unsigned int mask_sign_32 = {0x80000000, 0x80000000,
-                                              0x80000000, 0x80000000};
-  const __vector unsigned int mask_exp_32 = {0x7F800000, 0x7F800000, 0x7F800000,
-                                             0x7F800000};
-  const __vector unsigned int mask_mant_32 = {0x007FFFFF, 0x007FFFFF,
-                                              0x007FFFFF, 0x007FFFFF};
-
-  // Use SIGNED integers for exponent math to handle underflow check
-  const __vector signed int bias_adj = {112, 112, 112, 112};
-  const __vector signed int zero = {0, 0, 0, 0};
-  const __vector signed int max_exp = {31, 31, 31, 31};  // Max FP16 exp
-  const __vector unsigned int exp_max_fp32 = {0xFF, 0xFF, 0xFF, 0xFF};
-  const __vector unsigned int exp_max_fp16 = {0x1F, 0x1F, 0x1F, 0x1F};
-
-  __vector unsigned int s = (in & mask_sign_32) >> 16;
-  __vector unsigned int e_u = (in & mask_exp_32) >> 23;
-
-  // Check for NaN/Inf: exponent = 0xFF in FP32
-  __vector __bool int is_nan_inf = vec_cmpeq(e_u, exp_max_fp32);
-
-  __vector signed int e_s = (__vector signed int)e_u;
-  e_s = vec_sub(e_s, bias_adj);
-  e_s = vec_max(e_s, zero);
-  e_s = vec_min(e_s, max_exp);
-  __vector unsigned int e_normal = (__vector unsigned int)e_s;
-
-  __vector unsigned int e_final = vec_sel(e_normal, exp_max_fp16, is_nan_inf);
-
-  const __vector unsigned int one_v = {1, 1, 1, 1};
-  const __vector unsigned int mask_sticky = {0xFFF, 0xFFF, 0xFFF, 0xFFF};
-
-  __vector unsigned int round_bit = (in >> 12) & one_v;
-  __vector unsigned int sticky = in & mask_sticky;
-  __vector unsigned int m = (in & mask_mant_32) >> 13;
-  __vector unsigned int lsb = m & one_v;  // LSB of mantissa for tie-breaking
-
-  // Round up if: round_bit && (sticky || lsb)
-  __vector __bool int sticky_nonzero =
-      vec_cmpgt(sticky, (__vector unsigned int){0, 0, 0, 0});
-  __vector __bool int lsb_set = vec_cmpeq(lsb, one_v);
-  __vector __bool int round_up =
-      vec_and(vec_cmpeq(round_bit, one_v), vec_or(sticky_nonzero, lsb_set));
-
-  m = vec_sel(m, m + one_v, round_up);
-
-  const __vector unsigned int mant_mask = {0x3FF, 0x3FF, 0x3FF, 0x3FF};
-  const __vector unsigned int max_normal_exp = {0x1E, 0x1E, 0x1E, 0x1E};
-  __vector __bool int mant_overflows = vec_cmpgt(m, mant_mask);
-  __vector __bool int would_overflow_to_inf =
-      vec_and(mant_overflows, vec_cmpeq(e_final, max_normal_exp));
-  __vector unsigned int e_inc = vec_min(e_final + one_v, exp_max_fp16);
-  e_final = vec_sel(e_final, e_inc, mant_overflows);
-  m = vec_and(m, mant_mask);
-  e_final = vec_sel(e_final, max_normal_exp, would_overflow_to_inf);
-  m = vec_sel(m, mant_mask, would_overflow_to_inf);
-
-  return s | (e_final << 10) | m;
-}
-
 struct BF16Vec32 : public Vec<BF16Vec32> {
   constexpr static int VEC_ELEM_NUM = 32;
 
@@ -301,18 +180,6 @@ struct FP32Vec8 : public Vec<FP32Vec8> {
     reg.val[1] = (__vector float)vec_mergel(v.reg, zero);
   }
 
-  explicit FP32Vec8(const FP16Vec8& v) {
-    // Cast to UNSIGNED short vector to prevent sign-extension during unpack
-    __vector unsigned short raw_u = (__vector unsigned short)v.reg;
-
-    // Unpack 8x16-bit to two 4x32-bit vectors (Zero extended)
-    __vector unsigned int raw_hi = (__vector unsigned int)vec_unpackh(raw_u);
-    __vector unsigned int raw_lo = (__vector unsigned int)vec_unpackl(raw_u);
-
-    reg.val[0] = fp16_to_fp32_bits(raw_hi);
-    reg.val[1] = fp16_to_fp32_bits(raw_lo);
-  }
-
   float reduce_sum() const {
     AliasReg ar;
     ar.reg = reg;
@@ -664,22 +531,6 @@ struct FP32Vec16 : public Vec<FP32Vec16> {
     reg.val[3] = (__vector float)vec_mergel(v.reg.val[1], zero);
   }
 
-  explicit FP32Vec16(const FP16Vec16& v) {
-    __vector unsigned int raw_hi_0 =
-        (__vector unsigned int)vec_unpackh(v.reg.val[0]);
-    __vector unsigned int raw_lo_0 =
-        (__vector unsigned int)vec_unpackl(v.reg.val[0]);
-    reg.val[0] = fp16_to_fp32_bits(raw_hi_0);
-    reg.val[1] = fp16_to_fp32_bits(raw_lo_0);
-
-    __vector unsigned int raw_hi_1 =
-        (__vector unsigned int)vec_unpackh(v.reg.val[1]);
-    __vector unsigned int raw_lo_1 =
-        (__vector unsigned int)vec_unpackl(v.reg.val[1]);
-    reg.val[2] = fp16_to_fp32_bits(raw_hi_1);
-    reg.val[3] = fp16_to_fp32_bits(raw_lo_1);
-  }
-
   explicit FP32Vec16(const BF16Vec8& v) : FP32Vec16(FP32Vec8(v)) {}
 
   FP32Vec16 operator*(const FP32Vec16& b) const {
@@ -777,10 +628,8 @@ struct VecType<c10::BFloat16> {
   using vec_type = BF16Vec8;
 };
 
-template <>
-struct VecType<c10::Half> {
-  using vec_type = FP16Vec8;
-};
+// On s390x, FP16 (Half) is not natively supported, use FP32 vectors instead
+using FP16Vec16 = FP32Vec16;
 
 template <typename T>
 void storeFP32(float v, T* ptr) {
@@ -801,52 +650,6 @@ inline void storeFP32<c10::BFloat16>(float v, c10::BFloat16* ptr) {
   *ptr = *(v_ptr + 1);
 }
 
-template <>
-inline void storeFP32<::c10::Half>(float v, ::c10::Half* ptr) {
-  // Use bit-manipulation for IEEE FP32 to FP16 conversion since vector
-  // intrinsics for FP32 to FP16 conversion does not use IEEE rounding and can
-  // produce incorrect results for some inputs. Process each of the 4 vectors
-  // separately.
-  uint32_t in;
-  std::memcpy(&in, &v, sizeof(in));
-
-  uint32_t s = (in & 0x80000000) >> 16;  // Sign
-  uint32_t e = (in & 0x7F800000) >> 23;  // Exponent
-  uint32_t round_bit = (in >> 12) & 1;
-  uint32_t sticky = (in & 0xFFF) != 0;  // Any bits in [11..0]
-  uint32_t m = (in & 0x007FFFFF) >> 13;
-  uint32_t lsb = m & 1;  // LSB of mantissa for tie-breaking
-
-  // Check for NaN/Inf before rounding
-  bool is_nan_inf = (e == 0xFF);
-
-  if (round_bit && (sticky || lsb)) {
-    m++;
-    // Handle mantissa overflow: if m overflows 10 bits, increment exponent
-    if (m > 0x3FF) {
-      m = 0;
-      e++;
-    }
-  }
-
-  if (is_nan_inf) {
-    // NaN/Inf: preserve it
-    e = 0x1F;
-  } else {
-    // Normal: adjust bias (127 - 15), flush subnormals to zero
-    e = (e >= 112) ? (e - 112) : 0;
-    // If exponent overflows to Inf range, saturate to max normal FP16 value
-    if (e > 0x1E) {
-      e = 0x1E;   // Max normal exponent
-      m = 0x3FF;  // Max mantissa
-    }
-  }
-
-  uint16_t fp16 = (uint16_t)(s | (e << 10) | m);
-
-  *reinterpret_cast<uint16_t*>(ptr) = fp16;
-}
-
 #ifndef __VEC_CLASS_FP_NAN
   #define __VEC_CLASS_FP_NAN (1 << 6)
 #endif
@@ -1000,44 +803,6 @@ inline BF16Vec16::BF16Vec16(const FP32Vec16& v) {
   reg.val[1] = (__vector signed short)vec_perm(inp2, inp3, omask);
 }
 
-inline FP16Vec8::FP16Vec8(const FP32Vec8& v) {
-  // Use bit-manipulation for IEEE FP32 to FP16 conversion since vector
-  // intrinsics for FP32 to FP16 conversion does not use IEEE rounding and can
-  // produce incorrect results for some inputs. Process each of the 4 vectors
-  // separately.
-  __vector unsigned int res_hi = fp32_to_fp16_bits(v.reg.val[0]);
-  __vector unsigned int res_lo = fp32_to_fp16_bits(v.reg.val[1]);
-
-  const __vector unsigned char perm_pack = {
-      2,  3,  6,  7,  10, 11, 14, 15,  // Select lower 2 bytes from res_hi
-      18, 19, 22, 23, 26, 27, 30, 31   // Select lower 2 bytes from res_lo
-  };
-
-  reg = vec_perm((__vector signed short)res_hi, (__vector signed short)res_lo,
-                 perm_pack);
-}
-
-inline FP16Vec16::FP16Vec16(const FP32Vec16& v) {
-  // Use bit-manipulation for IEEE FP32 to FP16 conversion since vector
-  // intrinsics for FP32 to FP16 conversion does not use IEEE rounding and can
-  // produce incorrect results for some inputs. Process each of the 4 vectors
-  // separately.
-  __vector unsigned int res_0 = fp32_to_fp16_bits(v.reg.val[0]);
-  __vector unsigned int res_1 = fp32_to_fp16_bits(v.reg.val[1]);
-  __vector unsigned int res_2 = fp32_to_fp16_bits(v.reg.val[2]);
-  __vector unsigned int res_3 = fp32_to_fp16_bits(v.reg.val[3]);
-
-  const __vector unsigned char perm_pack = {
-      2,  3,  6,  7,  10, 11, 14, 15,  // Lower 2 bytes from first vector
-      18, 19, 22, 23, 26, 27, 30, 31   // Lower 2 bytes from second vector
-  };
-
-  reg.val[0] = vec_perm((__vector signed short)res_0,
-                        (__vector signed short)res_1, perm_pack);
-  reg.val[1] = vec_perm((__vector signed short)res_2,
-                        (__vector signed short)res_3, perm_pack);
-}
-
 // 1D softmax over `n` elements in `input`, writes result to `output`.
 // Uses FP32Vec8 for main body, scalar tail handling.
 // Requirement: n > 0

csrc/cpu/mla_decode.cpp

@@ -18,8 +18,8 @@ struct KernelVecType<float> {
 
 template <>
 struct KernelVecType<c10::Half> {
-#if defined(__powerpc64__)
-  // Power specific vector types
+#if defined(__powerpc64__) || defined(__s390x__)
+  // Power and s390x architecture-specific vector types
   using qk_load_vec_type = vec_op::FP32Vec16;
   using qk_vec_type = vec_op::FP32Vec16;
   using v_load_vec_type = vec_op::FP32Vec16;

docker/Dockerfile

@@ -506,6 +506,7 @@ RUN apt-get update -y \
         curl \
         sudo \
         python3-pip \
+        git \
         ffmpeg \
         libsm6 \
         libxext6 \

docker/Dockerfile.cpu

@@ -134,6 +134,7 @@ WORKDIR /vllm-workspace
 # Copy test requirements
 COPY requirements/test.in requirements/cpu-test.in
 
+# TODO: Update to 2.9.0 when there is a new build for intel_extension_for_pytorch for that version
 RUN \
     sed -i '/mamba_ssm/d' requirements/cpu-test.in && \
     remove_packages_not_supported_on_aarch64() { \

docs/design/io_processor_plugins.md

@@ -14,26 +14,8 @@ IOProcessorOutput = TypeVar("IOProcessorOutput")
 
 class IOProcessor(ABC, Generic[IOProcessorInput, IOProcessorOutput]):
     def __init__(self, vllm_config: VllmConfig):
-        super().__init__()
-
         self.vllm_config = vllm_config
 
-    @abstractmethod
-    def parse_data(self, data: object) -> IOProcessorInput:
-        raise NotImplementedError
-
-    def merge_sampling_params(
-        self,
-        params: SamplingParams | None = None,
-    ) -> SamplingParams:
-        return params or SamplingParams()
-
-    def merge_pooling_params(
-        self,
-        params: PoolingParams | None = None,
-    ) -> PoolingParams:
-        return params or PoolingParams()
-
     @abstractmethod
     def pre_process(
         self,
@@ -73,13 +55,29 @@ class IOProcessor(ABC, Generic[IOProcessorInput, IOProcessorOutput]):
             [(i, item) async for i, item in model_output], key=lambda output: output[0]
         )
         collected_output = [output[1] for output in sorted_output]
-        return self.post_process(collected_output, request_id=request_id, **kwargs)
+        return self.post_process(collected_output, request_id, **kwargs)
+
+    @abstractmethod
+    def parse_request(self, request: Any) -> IOProcessorInput:
+        raise NotImplementedError
+
+    def validate_or_generate_params(
+        self, params: SamplingParams | PoolingParams | None = None
+    ) -> SamplingParams | PoolingParams:
+        return params or PoolingParams()
+
+    @abstractmethod
+    def output_to_response(
+        self, plugin_output: IOProcessorOutput
+    ) -> IOProcessorResponse:
+        raise NotImplementedError
 ```
 
-The `parse_data` method is used for validating the user data and converting it into the input expected by the `pre_process*` methods.
-The `merge_sampling_params` and `merge_pooling_params` methods merge input `SamplingParams` or `PoolingParams` (if any) with the default one.
+The `parse_request` method is used for validating the user prompt and converting it into the input expected by the `pre_process`/`pre_process_async` methods.
 The `pre_process*` methods take the validated plugin input to generate vLLM's model prompts for regular inference.
 The `post_process*` methods take `PoolingRequestOutput` objects as input and generate a custom plugin output.
+The `validate_or_generate_params` method is used for validating with the plugin any `SamplingParameters`/`PoolingParameters` received with the user request, or to generate new ones if none are specified. The function always returns the validated/generated parameters.
+The `output_to_response` method is used only for online serving and converts the plugin output to the `IOProcessorResponse` type that is then returned by the API Server. The implementation of the `/pooling` serving endpoint is available here [vllm/entrypoints/openai/serving_pooling.py](../../vllm/entrypoints/pooling/pooling/serving.py).
 
 An example implementation of a plugin that enables generating geotiff images with the PrithviGeospatialMAE model is available [here](https://github.com/IBM/terratorch/tree/main/terratorch/vllm/plugins/segmentation). Please, also refer to our online ([examples/pooling/plugin/prithvi_geospatial_mae_online.py](../../examples/pooling/plugin/prithvi_geospatial_mae_online.py)) and offline ([examples/pooling/plugin/prithvi_geospatial_mae_io_processor.py](../../examples/pooling/plugin/prithvi_geospatial_mae_io_processor.py)) inference examples.
 

docs/getting_started/installation/gpu.xpu.inc.md

@@ -6,11 +6,10 @@ vLLM initially supports basic model inference and serving on Intel GPU platform.
 # --8<-- [start:requirements]
 
 - Supported Hardware: Intel Data Center GPU, Intel ARC GPU
-- OneAPI requirements: oneAPI 2025.3
-- Dependency: [vllm-xpu-kernels](https://github.com/vllm-project/vllm-xpu-kernels): a package provide all necessary vllm custom kernel when running vLLM on Intel GPU platform, 
+- OneAPI requirements: oneAPI 2025.1
 - Python: 3.12
 !!! warning
-    The provided vllm-xpu-kernels whl is Python3.12 specific so this version is a MUST.
+    The provided IPEX whl is Python3.12 specific so this version is a MUST.
 
 # --8<-- [end:requirements]
 # --8<-- [start:set-up-using-python]
@@ -25,7 +24,7 @@ Currently, there are no pre-built XPU wheels.
 # --8<-- [end:pre-built-wheels]
 # --8<-- [start:build-wheel-from-source]
 
-- First, install required [driver](https://dgpu-docs.intel.com/driver/installation.html#installing-gpu-drivers) and [Intel OneAPI](https://www.intel.com/content/www/us/en/developer/tools/oneapi/base-toolkit.html) 2025.3 or later.
+- First, install required [driver](https://dgpu-docs.intel.com/driver/installation.html#installing-gpu-drivers) and [Intel OneAPI](https://www.intel.com/content/www/us/en/developer/tools/oneapi/base-toolkit.html) 2025.1 or later.
 - Second, install Python packages for vLLM XPU backend building:
 
 ```bash
@@ -38,7 +37,7 @@ pip install -v -r requirements/xpu.txt
 - Then, build and install vLLM XPU backend:
 
 ```bash
-VLLM_TARGET_DEVICE=xpu pip install --no-build-isolation -e . -v
+VLLM_TARGET_DEVICE=xpu python setup.py install
 ```
 
 # --8<-- [end:build-wheel-from-source]

docs/models/supported_models.md

@@ -790,7 +790,6 @@ Speech2Text models trained specifically for Automatic Speech Recognition.
 
 | Architecture | Models | Example HF Models | [LoRA](../features/lora.md) | [PP](../serving/parallelism_scaling.md) |
 |--------------|--------|-------------------|----------------------|---------------------------|
-| `FunASRForConditionalGeneration` | FunASR | `allendou/Fun-ASR-Nano-2512-vllm`, etc. | | |
 | `Gemma3nForConditionalGeneration` | Gemma3n | `google/gemma-3n-E2B-it`, `google/gemma-3n-E4B-it`, etc. | | |
 | `GlmAsrForConditionalGeneration` | GLM-ASR | `zai-org/GLM-ASR-Nano-2512` | ✅︎ | ✅︎ |
 | `GraniteSpeechForConditionalGeneration` | Granite Speech | `ibm-granite/granite-speech-3.3-2b`, `ibm-granite/granite-speech-3.3-8b`, etc. | ✅︎ | ✅︎ |

examples/online_serving/openai_transcription_client.py

@@ -26,9 +26,7 @@ from openai import AsyncOpenAI, OpenAI
 from vllm.assets.audio import AudioAsset
 
 
-def sync_openai(
-    audio_path: str, client: OpenAI, model: str, *, repetition_penalty: float = 1.3
-):
+def sync_openai(audio_path: str, client: OpenAI, model: str):
     """
     Perform synchronous transcription using OpenAI-compatible API.
     """
@@ -42,7 +40,7 @@ def sync_openai(
             # Additional sampling params not provided by OpenAI API.
             extra_body=dict(
                 seed=4419,
-                repetition_penalty=repetition_penalty,
+                repetition_penalty=1.3,
             ),
         )
         print("transcription result [sync]:", transcription.text)
@@ -131,12 +129,7 @@ def main(args):
     print(f"Using model: {model}")
 
     # Run the synchronous function
-    sync_openai(
-        audio_path=args.audio_path if args.audio_path else mary_had_lamb,
-        client=client,
-        model=model,
-        repetition_penalty=args.repetition_penalty,
-    )
+    sync_openai(args.audio_path if args.audio_path else mary_had_lamb, client, model)
 
     # Run the asynchronous function
     if "openai" in model:
@@ -168,11 +161,5 @@ if __name__ == "__main__":
         default=None,
         help="The path to the audio file to transcribe.",
     )
-    parser.add_argument(
-        "--repetition_penalty",
-        type=float,
-        default=1.3,
-        help="repetition penalty",
-    )
     args = parser.parse_args()
     main(args)

requirements/common.txt

@@ -7,7 +7,7 @@ requests >= 2.26.0
 tqdm
 blake3
 py-cpuinfo
-transformers >= 4.56.0, < 5
+transformers @ git+https://github.com/huggingface/transformers.git@main
 tokenizers >= 0.21.1  # Required for fast incremental detokenization.
 protobuf # Required by LlamaTokenizer, gRPC.
 fastapi[standard] >= 0.115.0 # Required by FastAPI's form models in the OpenAI API server's audio transcriptions endpoint.

requirements/nightly_torch_test.txt

@@ -43,5 +43,5 @@ tritonclient>=2.51.0
 numba == 0.61.2 # Required for N-gram speculative decoding
 numpy
 runai-model-streamer[s3,gcs]==0.15.3
-fastsafetensors>=0.2.2
+fastsafetensors>=0.1.10
 pydantic>=2.12 # 2.11 leads to error on python 3.13

requirements/test.in

@@ -53,7 +53,7 @@ arctic-inference == 0.1.1 # Required for suffix decoding test
 numba == 0.61.2 # Required for N-gram speculative decoding
 numpy
 runai-model-streamer[s3,gcs]==0.15.3
-fastsafetensors>=0.2.2 # 0.2.2 contains important fixes for multi-GPU mem usage
+fastsafetensors>=0.1.10
 pydantic>=2.12 # 2.11 leads to error on python 3.13
 decord==0.6.0
 terratorch @ git+https://github.com/IBM/terratorch.git@1.1.rc3 # required for PrithviMAE test

requirements/test.txt

@@ -224,7 +224,7 @@ fastparquet==2024.11.0
     # via genai-perf
 fastrlock==0.8.2
     # via cupy-cuda12x
-fastsafetensors==0.2.2
+fastsafetensors==0.1.10
     # via -r requirements/test.in
 filelock==3.16.1
     # via
@@ -1174,6 +1174,7 @@ torch==2.10.0+cu129
     #   bitsandbytes
     #   efficientnet-pytorch
     #   encodec
+    #   fastsafetensors
     #   kornia
     #   lightly
     #   lightning

requirements/xpu.txt

@@ -15,4 +15,4 @@ torch==2.10.0+xpu
 torchaudio
 torchvision
 
-vllm_xpu_kernels @ https://github.com/vllm-project/vllm-xpu-kernels/releases/download/v0.1.2/vllm_xpu_kernels-0.1.2-cp312-cp312-linux_x86_64.whl
+vllm_xpu_kernels @ https://github.com/vllm-project/vllm-xpu-kernels/releases/download/v0.1.1/vllm_xpu_kernels-0.1.1-cp312-cp312-linux_x86_64.whl

setup.py

@@ -1035,7 +1035,7 @@ setup(
     extras_require={
         "bench": ["pandas", "matplotlib", "seaborn", "datasets", "scipy"],
         "tensorizer": ["tensorizer==2.10.1"],
-        "fastsafetensors": ["fastsafetensors >= 0.2.2"],
+        "fastsafetensors": ["fastsafetensors >= 0.1.10"],
         "runai": ["runai-model-streamer[s3,gcs] >= 0.15.3"],
         "audio": [
             "librosa",

tests/compile/fullgraph/test_simple.py

@@ -27,29 +27,10 @@ from ...utils import create_new_process_for_each_test
 from ..silly_attention import get_global_counter, reset_global_counter
 
 
-# Custom op that returns an unbacked symint during graph capture
-@torch.library.custom_op("mylib::foo", mutates_args=())
-def foo(x: torch.Tensor) -> int:
-    return 3
-
-
-@foo.register_fake
-def _(x):
-    return torch.library.get_ctx().new_dynamic_size()
-
-
 @support_torch_compile
 class SillyModel(nn.Module):
-    def __init__(
-        self,
-        *,
-        vllm_config: VllmConfig,
-        prefix: str = "",
-        intermediate_unbacked=False,
-        **kwargs,
-    ) -> None:
+    def __init__(self, *, vllm_config: VllmConfig, prefix: str = "", **kwargs) -> None:
         super().__init__()
-        self.intermediate_unbacked = intermediate_unbacked
 
     def forward(self, x: torch.Tensor) -> torch.Tensor:
         """
@@ -63,13 +44,6 @@ class SillyModel(nn.Module):
         torch.ops.silly.attention(x, x, x, out)
         x = out
         x = x - 2
-
-        if self.intermediate_unbacked:
-            # Test for unbacked symints: the following is a fancy way to multiply by 1
-            u0 = foo(x)
-            ones = x.new_ones(x.shape[0], u0).sum(-1) / 3
-            x = x * ones
-
         x = x - 1
         out = torch.empty_like(x)
         torch.ops.silly.attention(x, x, x, out)
@@ -78,7 +52,6 @@ class SillyModel(nn.Module):
         return x
 
 
-@torch._dynamo.config.patch(capture_dynamic_output_shape_ops=True)
 def _run_simple_model(
     splitting_ops,
     use_inductor_graph_partition,
@@ -87,8 +60,6 @@ def _run_simple_model(
     expected_num_piecewise_capturable_graphs_seen,
     expected_num_backend_compilations,
     expected_num_cudagraph_captured,
-    *,
-    intermediate_unbacked=False,
 ):
     vllm_config = VllmConfig(
         compilation_config=CompilationConfig(
@@ -101,11 +72,7 @@ def _run_simple_model(
         )
     )
     with set_current_vllm_config(vllm_config):
-        model = SillyModel(
-            vllm_config=vllm_config,
-            prefix="",
-            intermediate_unbacked=intermediate_unbacked,
-        )
+        model = SillyModel(vllm_config=vllm_config, prefix="")
 
     inputs = torch.randn(100).cuda()
 
@@ -158,10 +125,9 @@ def _run_simple_model(
 
 
 @pytest.mark.parametrize("backend", ["inductor", "eager"])
-@pytest.mark.parametrize("intermediate_unbacked", [True, False])
 @torch.inference_mode()
 @create_new_process_for_each_test("spawn")
-def test_simple_piecewise_compile(backend, intermediate_unbacked):
+def test_simple_piecewise_compile(backend):
     _run_simple_model(
         splitting_ops=["silly::attention"],
         use_inductor_graph_partition=False,
@@ -174,7 +140,6 @@ def test_simple_piecewise_compile(backend, intermediate_unbacked):
         expected_num_backend_compilations=3,
         # num_cudagraph_sizes * num_piecewise_capturable_graphs_seen
         expected_num_cudagraph_captured=6,
-        intermediate_unbacked=intermediate_unbacked,
     )
 
 

tests/kernels/quantization/test_rocm_skinny_gemms.py

@@ -270,9 +270,6 @@ def test_rocm_wvsplitk_fp8_kernel(
     out = ops.wvSplitKQ(B, A, dtype, scale_a, scale_b, get_cu_count(), BIAS)
 
     if xnorm:
-        torch.testing.assert_close(out, ref_out, atol=1e-3, rtol=1e-8)
-    elif k >= 32 * 1024:
-        # wider pytrch thresh for large-K & no xnorm
-        torch.testing.assert_close(out, ref_out, atol=0.07, rtol=5e-2)
+        assert torch.allclose(out, ref_out, atol=1e-3, rtol=1e-8)
     else:
-        torch.testing.assert_close(out, ref_out, atol=0.01, rtol=0.01)
+        assert torch.allclose(out, ref_out, 0.01)

tests/models/registry.py

@@ -713,10 +713,6 @@ _MULTIMODAL_EXAMPLE_MODELS = {
         "baidu/ERNIE-4.5-VL-28B-A3B-PT",
         trust_remote_code=True,
     ),
-    "FunASRForConditionalGeneration": _HfExamplesInfo(
-        "allendou/Fun-ASR-Nano-2512-vllm",
-        is_available_online=False,
-    ),
     "FunAudioChatForConditionalGeneration": _HfExamplesInfo(
         "funaudiochat", is_available_online=False
     ),

tests/plugins/prithvi_io_processor_plugin/prithvi_io_processor/prithvi_processor.py

@@ -18,10 +18,18 @@ from einops import rearrange
 from terratorch.datamodules import Sen1Floods11NonGeoDataModule
 
 from vllm.config import VllmConfig
+from vllm.entrypoints.pooling.pooling.protocol import (
+    IOProcessorRequest,
+    IOProcessorResponse,
+)
 from vllm.inputs.data import PromptType
 from vllm.logger import init_logger
 from vllm.outputs import PoolingRequestOutput
-from vllm.plugins.io_processors.interface import IOProcessor
+from vllm.plugins.io_processors.interface import (
+    IOProcessor,
+    IOProcessorInput,
+    IOProcessorOutput,
+)
 
 from .types import DataModuleConfig, ImagePrompt, ImageRequestOutput
 
@@ -219,7 +227,7 @@ def load_image(
     return imgs, temporal_coords, location_coords, metas
 
 
-class PrithviMultimodalDataProcessor(IOProcessor[ImagePrompt, ImageRequestOutput]):
+class PrithviMultimodalDataProcessor(IOProcessor):
     indices = [0, 1, 2, 3, 4, 5]
 
     def __init__(self, vllm_config: VllmConfig):
@@ -243,15 +251,34 @@ class PrithviMultimodalDataProcessor(IOProcessor[ImagePrompt, ImageRequestOutput
         self.requests_cache: dict[str, dict[str, Any]] = {}
         self.indices = DEFAULT_INPUT_INDICES
 
-    def parse_data(self, data: object) -> ImagePrompt:
-        if isinstance(data, dict):
-            return ImagePrompt(**data)
+    def parse_request(self, request: Any) -> IOProcessorInput:
+        if type(request) is dict:
+            image_prompt = ImagePrompt(**request)
+            return image_prompt
+        if isinstance(request, IOProcessorRequest):
+            if not hasattr(request, "data"):
+                raise ValueError("missing 'data' field in OpenAIBaseModel Request")
 
-        raise ValueError("Prompt data should be an `ImagePrompt`")
+            request_data = request.data
+
+            if type(request_data) is dict:
+                return ImagePrompt(**request_data)
+            else:
+                raise ValueError("Unable to parse the request data")
+
+        raise ValueError("Unable to parse request")
+
+    def output_to_response(
+        self, plugin_output: IOProcessorOutput
+    ) -> IOProcessorResponse:
+        return IOProcessorResponse(
+            request_id=plugin_output.request_id,
+            data=plugin_output,
+        )
 
     def pre_process(
         self,
-        prompt: ImagePrompt,
+        prompt: IOProcessorInput,
         request_id: str | None = None,
         **kwargs,
     ) -> PromptType | Sequence[PromptType]:
@@ -337,7 +364,7 @@ class PrithviMultimodalDataProcessor(IOProcessor[ImagePrompt, ImageRequestOutput
         model_output: Sequence[PoolingRequestOutput],
         request_id: str | None = None,
         **kwargs,
-    ) -> ImageRequestOutput:
+    ) -> IOProcessorOutput:
         pred_imgs_list = []
 
         if request_id and (request_id in self.requests_cache):
@@ -382,7 +409,5 @@ class PrithviMultimodalDataProcessor(IOProcessor[ImagePrompt, ImageRequestOutput
         )
 
         return ImageRequestOutput(
-            type=out_format,
-            format="tiff",
-            data=out_data,
+            type=out_format, format="tiff", data=out_data, request_id=request_id
         )

tests/plugins/prithvi_io_processor_plugin/prithvi_io_processor/types.py

@@ -38,6 +38,9 @@ class ImagePrompt(BaseModel):
     """
 
 
+MultiModalPromptType = ImagePrompt
+
+
 class ImageRequestOutput(BaseModel):
     """
     The output data of an image request to vLLM.
@@ -51,3 +54,4 @@ class ImageRequestOutput(BaseModel):
     type: Literal["path", "b64_json"]
     format: str
     data: str
+    request_id: str | None = None

tests/plugins_tests/test_io_processor_plugins.py

@@ -75,7 +75,9 @@ async def test_prithvi_mae_plugin_online(
     # verify the output is formatted as expected for this plugin
     plugin_data = parsed_response.data
 
-    assert all(plugin_data.get(attr) for attr in ["type", "format", "data"])
+    assert all(
+        plugin_data.get(attr) for attr in ["type", "format", "data", "request_id"]
+    )
 
     # We just check that the output is a valid base64 string.
     # Raises an exception and fails the test if the string is corrupted.
@@ -108,7 +110,9 @@ def test_prithvi_mae_plugin_offline(vllm_runner, model_name: str):
     output = pooler_output[0].outputs
 
     # verify the output is formatted as expected for this plugin
-    assert all(hasattr(output, attr) for attr in ["type", "format", "data"])
+    assert all(
+        hasattr(output, attr) for attr in ["type", "format", "data", "request_id"]
+    )
 
     # We just check that the output is a valid base64 string.
     # Raises an exception and fails the test if the string is corrupted.

tests/quantization/test_cpu_wna16.py

@@ -17,7 +17,7 @@ DTYPE = ["bfloat16"]
 
 @pytest.mark.parametrize("model", MODELS)
 @pytest.mark.parametrize("dtype", DTYPE)
-def test_cpu_quant(vllm_runner, model, dtype):
+def test_ipex_quant(vllm_runner, model, dtype):
     with vllm_runner(model, dtype=dtype) as llm:
         output = llm.generate_greedy(["The capital of France is"], max_tokens=32)
     assert output

tests/quantization/test_ipex_quant.py

@@ -0,0 +1,32 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+"""Test model set-up and inference for quantized HF models supported
+on the CPU/GPU backend using IPEX (including AWQ/GPTQ).
+
+Validating the configuration and printing results for manual checking.
+
+Run `pytest tests/quantization/test_ipex_quant.py`.
+"""
+
+import pytest
+
+from vllm.platforms import current_platform
+
+MODELS = [
+    "AMead10/Llama-3.2-1B-Instruct-AWQ",
+    "shuyuej/Llama-3.2-1B-Instruct-GPTQ",  # with g_idx
+]
+DTYPE = ["bfloat16"]
+
+
+@pytest.mark.skipif(
+    not current_platform.is_cpu() and not current_platform.is_xpu(),
+    reason="only supports Intel CPU/XPU backend.",
+)
+@pytest.mark.parametrize("model", MODELS)
+@pytest.mark.parametrize("dtype", DTYPE)
+def test_ipex_quant(vllm_runner, model, dtype):
+    with vllm_runner(model, dtype=dtype, enforce_eager=True) as llm:
+        output = llm.generate_greedy(["The capital of France is"], max_tokens=4)
+    assert output
+    print(output)

tests/v1/engine/test_process_multi_modal_uuids.py

@@ -20,6 +20,7 @@ def _build_input_processor(
 ) -> InputProcessor:
     model_config = ModelConfig(
         model="Qwen/Qwen2.5-VL-3B-Instruct",
+        skip_tokenizer_init=True,
         max_model_len=128,
         mm_processor_cache_gb=mm_cache_gb,
     )

vllm/_ipex_ops.py

@@ -53,7 +53,7 @@ if hasattr(torch.ops._xpu_C, "int4_gemm_w4a16"):
         return torch.empty((M, N), dtype=input.dtype, device=input.device)
 
 
-class xpu_ops:
+class ipex_ops:
     @staticmethod
     def flash_attn_varlen_func(
         q: torch.Tensor,
@@ -73,7 +73,7 @@ class xpu_ops:
         cu_seqlens_k: torch.Tensor | None = None,
         # passed in qwen vl
         dropout_p: float = 0.0,
-        # The following parameters are not used in xpu kernel currently,
+        # The following parameters are not used in ipex kernel currently,
         # we keep API compatible to CUDA's.
         scheduler_metadata=None,
         fa_version: int = 2,
@@ -153,6 +153,6 @@ class xpu_ops:
         sm_margin=0,  # Can be tuned if some SMs are used for communication
     ) -> None:
         logger.warning_once(
-            "get_scheduler_metadata is not implemented for xpu_ops, returning None."
+            "get_scheduler_metadata is not implemented for ipex_ops, returning None."
         )
         return None

vllm/benchmarks/datasets.py

@@ -1310,11 +1310,6 @@ class _ValidateDatasetArgs(argparse.Action):
 
 
 def add_dataset_parser(parser: FlexibleArgumentParser):
-    parser.add_argument(
-        "--trust-remote-code",
-        action="store_true",
-        help="Trust remote code from huggingface",
-    )
     parser.add_argument("--seed", type=int, default=0)
     parser.add_argument(
         "--num-prompts",

vllm/benchmarks/serve.py

@@ -1313,6 +1313,11 @@ def add_cli_args(parser: argparse.ArgumentParser):
         "bursty requests. A higher burstiness value (burstiness > 1) "
         "results in a more uniform arrival of requests.",
     )
+    parser.add_argument(
+        "--trust-remote-code",
+        action="store_true",
+        help="Trust remote code from huggingface",
+    )
     parser.add_argument(
         "--disable-tqdm",
         action="store_true",

vllm/compilation/backends.py

@@ -570,9 +570,7 @@ class PiecewiseCompileInterpreter(torch.fx.Interpreter):  # type: ignore[misc]
     ) -> Any:
         assert isinstance(target, str)
 
-        gm = getattr(self.module, target)
-        outputs = gm.graph.output_node().args[0]
-        output = fx.map_arg(outputs, lambda node: node.meta["example_value"])
+        output = super().call_module(target, args, kwargs)
 
         if target in self.compile_submod_names:
             index = self.compile_submod_names.index(target)

vllm/config/compilation.py

@@ -115,7 +115,7 @@ class PassConfig:
     """Fuse the custom SiluMul + quant ops."""
     fuse_attn_quant: bool = Field(default=None)
     """Fuse the custom attention + quant ops."""
-    eliminate_noops: bool = Field(default=True)
+    eliminate_noops: bool = Field(default=None)
     """Eliminate no-op ops."""
     enable_sp: bool = Field(default=None)
     """Enable sequence parallelism."""
@@ -194,6 +194,7 @@ class PassConfig:
         "fuse_norm_quant",
         "fuse_act_quant",
         "fuse_attn_quant",
+        "eliminate_noops",
         "enable_sp",
         "fuse_gemm_comms",
         "fuse_allreduce_rms",

vllm/config/vllm.py

@@ -102,19 +102,6 @@ def enable_act_fusion(cfg: "VllmConfig") -> bool:
     ) or cfg.compilation_config.is_custom_op_enabled("quant_fp8")
 
 
-def enable_allreduce_rms_fusion(cfg: "VllmConfig") -> bool:
-    """Enable if TP > 1 and Hopper+ and flashinfer installed."""
-    from vllm.platforms import current_platform
-    from vllm.utils.flashinfer import has_flashinfer
-
-    return (
-        cfg.parallel_config.tensor_parallel_size > 1
-        and current_platform.is_cuda()
-        and current_platform.has_device_capability(90)
-        and has_flashinfer()
-    )
-
-
 def enable_norm_pad_fusion(cfg: "VllmConfig") -> bool:
     """Enable if using AITER RMSNorm and AITER Triton GEMMs
     and hidden size is 2880 i.e. gpt-oss; otherwise Inductor handles fusion."""
@@ -131,6 +118,7 @@ def enable_norm_pad_fusion(cfg: "VllmConfig") -> bool:
 OPTIMIZATION_LEVEL_00 = {
     "compilation_config": {
         "pass_config": {
+            "eliminate_noops": False,
             "fuse_norm_quant": False,
             "fuse_act_quant": False,
             "fuse_allreduce_rms": False,
@@ -149,6 +137,7 @@ OPTIMIZATION_LEVEL_00 = {
 OPTIMIZATION_LEVEL_01 = {
     "compilation_config": {
         "pass_config": {
+            "eliminate_noops": True,
             "fuse_norm_quant": enable_norm_fusion,
             "fuse_act_quant": enable_act_fusion,
             "fuse_allreduce_rms": False,
@@ -167,9 +156,10 @@ OPTIMIZATION_LEVEL_01 = {
 OPTIMIZATION_LEVEL_02 = {
     "compilation_config": {
         "pass_config": {
+            "eliminate_noops": True,
             "fuse_norm_quant": enable_norm_fusion,
             "fuse_act_quant": enable_act_fusion,
-            "fuse_allreduce_rms": enable_allreduce_rms_fusion,
+            "fuse_allreduce_rms": False,
             "fuse_attn_quant": IS_QUANTIZED,
             "enable_sp": IS_DENSE,
             "fuse_gemm_comms": IS_DENSE,
@@ -185,9 +175,10 @@ OPTIMIZATION_LEVEL_02 = {
 OPTIMIZATION_LEVEL_03 = {
     "compilation_config": {
         "pass_config": {
+            "eliminate_noops": True,
             "fuse_norm_quant": enable_norm_fusion,
             "fuse_act_quant": enable_act_fusion,
-            "fuse_allreduce_rms": enable_allreduce_rms_fusion,
+            "fuse_allreduce_rms": False,
             "fuse_attn_quant": IS_QUANTIZED,
             "enable_sp": IS_DENSE,
             "fuse_gemm_comms": IS_DENSE,

vllm/entrypoints/llm.py

@@ -85,6 +85,7 @@ from vllm.tasks import PoolingTask
 from vllm.tokenizers import TokenizerLike
 from vllm.tokenizers.mistral import MistralTokenizer
 from vllm.usage.usage_lib import UsageContext
+from vllm.utils.collection_utils import as_iter, is_list_of
 from vllm.utils.counter import Counter
 from vllm.v1.engine.llm_engine import LLMEngine
 from vllm.v1.sample.logits_processor import LogitsProcessor
@@ -94,7 +95,6 @@ if TYPE_CHECKING:
 
 logger = init_logger(__name__)
 
-_P = TypeVar("_P", bound=SamplingParams | PoolingParams | None)
 _R = TypeVar("_R", default=Any)
 
 
@@ -1056,7 +1056,9 @@ class LLM:
                 dict(truncate_prompt_tokens=truncate_prompt_tokens),
             )
 
-        if use_io_processor := (isinstance(prompts, dict) and "data" in prompts):
+        io_processor_prompt = False
+        if isinstance(prompts, dict) and "data" in prompts:
+            io_processor_prompt = True
             if self.io_processor is None:
                 raise ValueError(
                     "No IOProcessor plugin installed. Please refer "
@@ -1066,42 +1068,40 @@ class LLM:
                 )
 
             # Validate the request data is valid for the loaded plugin
-            validated_prompt = self.io_processor.parse_data(prompts)
+            validated_prompt = self.io_processor.parse_request(prompts)
 
             # obtain the actual model prompts from the pre-processor
             prompts = self.io_processor.pre_process(prompt=validated_prompt)
-            prompts_seq = prompt_to_seq(prompts)
 
-            params_seq: Sequence[PoolingParams] = [
-                self.io_processor.merge_pooling_params(param)
-                for param in self._params_to_seq(
-                    pooling_params,
-                    len(prompts_seq),
-                )
-            ]
-            for p in params_seq:
-                if p.task is None:
-                    p.task = "plugin"
-        else:
-            if pooling_params is None:
-                # Use default pooling params.
-                pooling_params = PoolingParams()
-
-            prompts_seq = prompt_to_seq(prompts)
-            params_seq = self._params_to_seq(pooling_params, len(prompts_seq))
-
-            for param in params_seq:
-                if param.task is None:
-                    param.task = pooling_task
-                elif param.task != pooling_task:
-                    msg = (
-                        f"You cannot overwrite {param.task=!r} with {pooling_task=!r}!"
+        if io_processor_prompt:
+            assert self.io_processor is not None
+            if is_list_of(pooling_params, PoolingParams):
+                validated_pooling_params: list[PoolingParams] = []
+                for param in as_iter(pooling_params):
+                    validated_pooling_params.append(
+                        self.io_processor.validate_or_generate_params(param)
                     )
-                    raise ValueError(msg)
+                pooling_params = validated_pooling_params
+            else:
+                assert not isinstance(pooling_params, Sequence)
+                pooling_params = self.io_processor.validate_or_generate_params(
+                    pooling_params
+                )
+
+        if pooling_params is None:
+            # Use default pooling params.
+            pooling_params = PoolingParams()
+
+        for param in as_iter(pooling_params):
+            if param.task is None:
+                param.task = pooling_task
+            elif param.task != pooling_task:
+                msg = f"You cannot overwrite {param.task=!r} with {pooling_task=!r}!"
+                raise ValueError(msg)
 
         outputs = self._run_completion(
-            prompts=prompts_seq,
-            params=params_seq,
+            prompts=prompts,
+            params=pooling_params,
             use_tqdm=use_tqdm,
             lora_request=lora_request,
             tokenization_kwargs=tokenization_kwargs,
@@ -1111,10 +1111,12 @@ class LLM:
             outputs, PoolingRequestOutput
         )
 
-        if use_io_processor:
+        if io_processor_prompt:
             # get the post-processed model outputs
             assert self.io_processor is not None
-            processed_outputs = self.io_processor.post_process(model_outputs)
+            processed_outputs = self.io_processor.post_process(
+                model_output=model_outputs
+            )
 
             return [
                 PoolingRequestOutput[Any](
@@ -1660,9 +1662,11 @@ class LLM:
 
     def _params_to_seq(
         self,
-        params: _P | Sequence[_P],
+        params: SamplingParams
+        | PoolingParams
+        | Sequence[SamplingParams | PoolingParams],
         num_requests: int,
-    ) -> Sequence[_P]:
+    ) -> Sequence[SamplingParams | PoolingParams]:
         if isinstance(params, Sequence):
             if len(params) != num_requests:
                 raise ValueError(

vllm/entrypoints/pooling/pooling/protocol.py

@@ -100,6 +100,9 @@ class IOProcessorRequest(PoolingBasicRequestMixin, EncodingRequestMixin, Generic
     data: T
     task: PoolingTask = "plugin"
 
+    def to_pooling_params(self):
+        return PoolingParams(task=self.task)
+
 
 class IOProcessorResponse(OpenAIBaseModel, Generic[T]):
     request_id: str | None = None

vllm/entrypoints/pooling/pooling/serving.py

@@ -85,6 +85,7 @@ class OpenAIServingPooling(OpenAIServing):
         request_id = f"pool-{self._base_request_id(raw_request)}"
         created_time = int(time.time())
 
+        is_io_processor_request = isinstance(request, IOProcessorRequest)
         try:
             lora_request = self._maybe_get_adapters(request)
 
@@ -94,7 +95,7 @@ class OpenAIServingPooling(OpenAIServing):
                 )
 
             engine_prompts: Sequence[PromptType | TokPrompt]
-            if use_io_processor := isinstance(request, IOProcessorRequest):
+            if is_io_processor_request:
                 if self.io_processor is None:
                     raise ValueError(
                         "No IOProcessor plugin installed. Please refer "
@@ -103,7 +104,7 @@ class OpenAIServingPooling(OpenAIServing):
                         "offline inference example for more details."
                     )
 
-                validated_prompt = self.io_processor.parse_data(request.data)
+                validated_prompt = self.io_processor.parse_request(request)
 
                 raw_prompts = await self.io_processor.pre_process_async(
                     prompt=validated_prompt, request_id=request_id
@@ -140,18 +141,13 @@ class OpenAIServingPooling(OpenAIServing):
         # Schedule the request and get the result generator.
         generators: list[AsyncGenerator[PoolingRequestOutput, None]] = []
         try:
-            if use_io_processor:
-                assert self.io_processor is not None
-
-                pooling_params = self.io_processor.merge_pooling_params()
-                if pooling_params.task is None:
-                    pooling_params.task = "plugin"
-
-                tokenization_kwargs: dict[str, Any] = {}
+            if is_io_processor_request:
+                assert self.io_processor is not None and isinstance(
+                    request, IOProcessorRequest
+                )
+                pooling_params = self.io_processor.validate_or_generate_params()
             else:
-                pooling_params = request.to_pooling_params()  # type: ignore
-                tok_params = request.build_tok_params(self.model_config)  # type: ignore
-                tokenization_kwargs = tok_params.get_encode_kwargs()
+                pooling_params = request.to_pooling_params()
 
             for i, engine_prompt in enumerate(engine_prompts):
                 request_id_item = f"{request_id}-{i}"
@@ -169,6 +165,12 @@ class OpenAIServingPooling(OpenAIServing):
                     else await self._get_trace_headers(raw_request.headers)
                 )
 
+                if is_io_processor_request:
+                    tokenization_kwargs: dict[str, Any] = {}
+                else:
+                    tok_params = request.build_tok_params(self.model_config)  # type: ignore
+                    tokenization_kwargs = tok_params.get_encode_kwargs()
+
                 generator = self.engine_client.encode(
                     engine_prompt,
                     pooling_params,
@@ -185,31 +187,13 @@ class OpenAIServingPooling(OpenAIServing):
 
         result_generator = merge_async_iterators(*generators)
 
-        if use_io_processor:
+        if is_io_processor_request:
             assert self.io_processor is not None
             output = await self.io_processor.post_process_async(
-                result_generator,
+                model_output=result_generator,
                 request_id=request_id,
             )
-
-            if callable(
-                output_to_response := getattr(
-                    self.io_processor, "output_to_response", None
-                )
-            ):
-                logger.warning_once(
-                    "`IOProcessor.output_to_response` is deprecated. To ensure "
-                    "consistency between offline and online APIs, "
-                    "`IOProcessorResponse` will become a transparent wrapper "
-                    "around output data from v0.19 onwards.",
-                )
-
-                if hasattr(output, "request_id") and output.request_id is None:
-                    output.request_id = request_id  # type: ignore
-
-                return output_to_response(output)  # type: ignore
-
-            return IOProcessorResponse(request_id=request_id, data=output)
+            return self.io_processor.output_to_response(output)
 
         assert isinstance(request, (PoolingCompletionRequest, PoolingChatRequest))
         num_prompts = len(engine_prompts)

vllm/model_executor/layers/fused_moe/__init__.py

@@ -102,7 +102,6 @@ if HAS_TRITON:
     )
     from vllm.model_executor.layers.fused_moe.xpu_fused_moe import (
         XPUExperts,
-        XPUExpertsFp8,
     )
 
     __all__ += [
@@ -122,7 +121,6 @@ if HAS_TRITON:
         "BatchedDeepGemmExperts",
         "TritonOrDeepGemmExperts",
         "XPUExperts",
-        "XPUExpertsFp8",
     ]
 else:
     # Some model classes directly use the custom ops. Add placeholders

vllm/model_executor/layers/fused_moe/fused_moe.py

@@ -95,19 +95,19 @@ def fused_moe_kernel_gptq_awq(
     # moving by 1 element in a particular dimension. E.g. `stride_am` is
     # how much to increase `a_ptr` by to get the element one row down
     # (A has M rows).
-    stride_am: tl.int64,
-    stride_ak: tl.int64,
-    stride_be: tl.int64,
-    stride_bk: tl.int64,
-    stride_bn: tl.int64,
-    stride_cm: tl.int64,
-    stride_cn: tl.int64,
-    stride_bse: tl.int64,
-    stride_bsk: tl.int64,
-    stride_bsn: tl.int64,
-    stride_bze: tl.int64,
-    stride_bzk: tl.int64,
-    stride_bzn: tl.int64,
+    stride_am,
+    stride_ak,
+    stride_be,
+    stride_bk,
+    stride_bn,
+    stride_cm,
+    stride_cn,
+    stride_bse,
+    stride_bsk,
+    stride_bsn,
+    stride_bze,
+    stride_bzk,
+    stride_bzn,
     block_k_diviable: tl.constexpr,
     group_size: tl.constexpr,
     # Meta-parameters
@@ -329,20 +329,20 @@ def fused_moe_kernel(
     # moving by 1 element in a particular dimension. E.g. `stride_am` is
     # how much to increase `a_ptr` by to get the element one row down
     # (A has M rows).
-    stride_am: tl.int64,
-    stride_ak: tl.int64,
-    stride_be: tl.int64,
-    stride_bk: tl.int64,
-    stride_bn: tl.int64,
-    stride_cm: tl.int64,
-    stride_cn: tl.int64,
-    stride_asm: tl.int64,
-    stride_ask: tl.int64,
-    stride_bse: tl.int64,
-    stride_bsk: tl.int64,
-    stride_bsn: tl.int64,
-    stride_bbe: tl.int64,  # bias expert stride
-    stride_bbn: tl.int64,  # bias N stride
+    stride_am,
+    stride_ak,
+    stride_be,
+    stride_bk,
+    stride_bn,
+    stride_cm,
+    stride_cn,
+    stride_asm,
+    stride_ask,
+    stride_bse,
+    stride_bsk,
+    stride_bsn,
+    stride_bbe,  # bias expert stride
+    stride_bbn,  # bias N stride
     # Block size for block-wise quantization
     group_n: tl.constexpr,
     group_k: tl.constexpr,

vllm/model_executor/layers/fused_moe/oracle/fp8.py

@@ -52,7 +52,6 @@ class Fp8MoeBackend(Enum):
     AITER = "AITER"
     VLLM_CUTLASS = "VLLM_CUTLASS"
     BATCHED_VLLM_CUTLASS = "BATCHED_VLLM_CUTLASS"
-    XPU = "XPU"
 
 
 def backend_to_kernel_cls(
@@ -124,13 +123,6 @@ def backend_to_kernel_cls(
 
         return CutlassBatchedExpertsFp8
 
-    elif backend == Fp8MoeBackend.XPU:
-        from vllm.model_executor.layers.fused_moe.xpu_fused_moe import (
-            XPUExpertsFp8,
-        )
-
-        return XPUExpertsFp8
-
     else:
         raise ValueError(f"Unknown FP8 MoE backend: {backend.value}")
 
@@ -162,7 +154,6 @@ def select_fp8_moe_backend(
         Fp8MoeBackend.TRITON,
         Fp8MoeBackend.BATCHED_TRITON,
         Fp8MoeBackend.MARLIN,
-        Fp8MoeBackend.XPU,
     ]
 
     # NOTE(rob): We need to peak into the P/F selection to determine
@@ -402,7 +393,6 @@ def convert_to_fp8_moe_kernel_format(
             Fp8MoeBackend.BATCHED_TRITON,
             Fp8MoeBackend.VLLM_CUTLASS,
             Fp8MoeBackend.BATCHED_VLLM_CUTLASS,
-            Fp8MoeBackend.XPU,
         ]:
             raise ValueError(f"Unsupported FP8 MoE backend: {fp8_backend.value}")
 

vllm/model_executor/layers/fused_moe/xpu_fused_moe.py

@@ -4,16 +4,13 @@ import torch
 
 import vllm.model_executor.layers.fused_moe.modular_kernel as mk
 from vllm.model_executor.layers.fused_moe.config import (
-    FusedMoEConfig,
     FusedMoEParallelConfig,
-    FusedMoEQuantConfig,
 )
 from vllm.model_executor.layers.fused_moe.topk_weight_and_reduce import (
     TopKWeightAndReduceNoOP,
 )
 from vllm.model_executor.layers.quantization.utils.quant_utils import (
     QuantKey,
-    kFp8DynamicTensorSym,
     kFp8StaticTensorSym,
 )
 from vllm.platforms import current_platform
@@ -23,21 +20,6 @@ if current_platform.is_xpu():
 
 
 class XPUExperts(mk.FusedMoEPermuteExpertsUnpermute):
-    def __init__(
-        self,
-        moe_config: FusedMoEConfig,
-        quant_config: FusedMoEQuantConfig,
-        max_num_tokens: int | None = None,
-        num_dispatchers: int | None = None,
-    ):
-        super().__init__(
-            moe_config,
-            quant_config,
-            max_num_tokens,
-            num_dispatchers,
-        )
-        self.is_fp8 = False
-
     @property
     def expects_unquantized_inputs(self) -> bool:
         return True
@@ -67,10 +49,10 @@ class XPUExperts(mk.FusedMoEPermuteExpertsUnpermute):
         weight_key: QuantKey | None,
         activation_key: QuantKey | None,
     ) -> bool:
+        # TODO: dispatch based on device.
         SUPPORTED_W_A = [
             (None, None),
             (kFp8StaticTensorSym, None),
-            (kFp8StaticTensorSym, kFp8DynamicTensorSym),
         ]
         return (weight_key, activation_key) in SUPPORTED_W_A
 
@@ -121,10 +103,10 @@ class XPUExperts(mk.FusedMoEPermuteExpertsUnpermute):
         xpu_fused_moe(
             hidden_states=hidden_states,
             w13=w1,
-            w13_scales=self.w1_scale,
+            w13_scales=a1q_scale,
             w13_bias=self.w1_bias,
             w2=w2,
-            w2_scales=self.w2_scale,
+            w2_scales=a2_scale,
             w2_bias=self.w2_bias,
             topk_weights=topk_weights,
             topk_ids=topk_ids,
@@ -134,22 +116,5 @@ class XPUExperts(mk.FusedMoEPermuteExpertsUnpermute):
             ep_rank=self.moe_config.ep_rank,
             ep_size=self.moe_config.ep_size,
             output=output,
-            is_fp8=self.is_fp8,
         )
-
-
-class XPUExpertsFp8(XPUExperts):
-    def __init__(
-        self,
-        moe_config: FusedMoEConfig,
-        quant_config: FusedMoEQuantConfig,
-        max_num_tokens: int | None = None,
-        num_dispatchers: int | None = None,
-    ):
-        super().__init__(
-            moe_config,
-            quant_config,
-            max_num_tokens,
-            num_dispatchers,
-        )
-        self.is_fp8 = True
+        return

vllm/model_executor/layers/quantization/mxfp4.py

@@ -160,7 +160,7 @@ def get_mxfp4_backend(with_lora_support: bool) -> Mxfp4Backend:
             logger.info_once("Using Triton backend")
             return Mxfp4Backend.TRITON
     elif current_platform.is_xpu():
-        logger.info_once("Using xpu backend on XPU")
+        logger.info_once("Using ipex marlin backend on XPU")
         return Mxfp4Backend.MARLIN
     elif current_platform.is_rocm() and has_triton_kernels():
         logger.info_once("Using Triton backend")

vllm/model_executor/layers/quantization/utils/flashinfer_fp4_moe.py

@@ -25,7 +25,6 @@ from vllm.model_executor.layers.quantization.utils.quant_utils import (
 from vllm.platforms import current_platform
 
 if TYPE_CHECKING:
-    from vllm.model_executor.layers.fused_moe.layer import FusedMoE
     from vllm.model_executor.layers.fused_moe.oracle.nvfp4 import (
         NvFp4MoeBackend,
     )
@@ -83,12 +82,8 @@ def _supports_routing_method(
 
 
 def _supports_parallel_config(moe_parallel_config: FusedMoEParallelConfig) -> bool:
-    """
-    TRTLLM is a monolithic kernel that requires dispatch_router_logits() for
-    the naive dispatch/combine path. DeepEP HT only implements dispatch() for
-    the modular kernel path, so TRTLLM is incompatible with DeepEP HT.
-    """
-    return not moe_parallel_config.use_deepep_ht_kernels
+    """Supports EP."""
+    return True
 
 
 def is_supported_config_trtllm(
@@ -317,7 +312,11 @@ def flashinfer_trtllm_fp4_moe(
     if use_llama4_routing:
         routing_method_type = flashinfer.RoutingMethodType.Llama4
 
-    # Cast to Fp32 (required by kernel).
+    # Prepare routing bias
+    routing_bias = e_score_correction_bias
+    if routing_bias is not None:
+        routing_bias = routing_bias.to(torch.bfloat16)
+
     router_logits = (
         router_logits.to(torch.float32)
         if routing_method_type == RoutingMethodType.DeepSeekV3
@@ -327,7 +326,7 @@ def flashinfer_trtllm_fp4_moe(
     # Call TRT-LLM FP4 block-scale MoE kernel
     out = flashinfer.fused_moe.trtllm_fp4_block_scale_moe(
         routing_logits=router_logits,
-        routing_bias=e_score_correction_bias,
+        routing_bias=routing_bias,
         hidden_states=hidden_states_fp4,
         hidden_states_scale=hidden_states_scale_linear_fp4.view(
             torch.float8_e4m3fn
@@ -444,7 +443,7 @@ def flashinfer_trtllm_fp4_routed_moe(
 
 def prepare_nvfp4_moe_layer_for_fi_or_cutlass(
     backend: "NvFp4MoeBackend",
-    layer: "FusedMoE",
+    layer: torch.nn.Module,
     w13: torch.Tensor,
     w13_scale: torch.Tensor,
     w13_scale_2: torch.Tensor,

vllm/model_executor/layers/sparse_attn_indexer.py

@@ -20,7 +20,7 @@ from vllm.v1.worker.workspace import current_workspace_manager
 if current_platform.is_cuda_alike():
     from vllm import _custom_ops as ops
 elif current_platform.is_xpu():
-    from vllm._xpu_ops import xpu_ops as ops
+    from vllm._ipex_ops import ipex_ops as ops
 
 logger = init_logger(__name__)
 

vllm/model_executor/models/qwen3_5.py

@@ -206,7 +206,7 @@ class Qwen3_5GatedDeltaNet(Qwen3NextGatedDeltaNet):
             output_size=self.num_v_heads,
             bias=False,
             quant_config=quant_config,
-            prefix=f"{prefix}.in_proj_b",
+            prefix=f"{prefix}.in_proj_ba",
         )
         self.in_proj_a = ColumnParallelLinear(
             input_size=self.hidden_size,

vllm/model_executor/models/registry.py

@@ -325,7 +325,6 @@ _MULTIMODAL_MODELS = {
         "ernie45_vl",
         "Ernie4_5_VLMoeForConditionalGeneration",
     ),
-    "FunASRForConditionalGeneration": ("funasr", "FunASRForConditionalGeneration"),  # noqa: E501
     "FunAudioChatForConditionalGeneration": (
         "funaudiochat",
         "FunAudioChatForConditionalGeneration",

vllm/multimodal/encoder_budget.py

@@ -62,7 +62,6 @@ class MultiModalBudget:
             processor = mm_registry.create_processor(model_config, cache=cache)
 
             self.cache = cache
-            self.processor = processor
             mm_config = model_config.get_multimodal_config()
             enable_mm_embeds = mm_config is not None and mm_config.enable_mm_embeds
 

vllm/platforms/cpu.py

@@ -345,6 +345,7 @@ class CpuPlatform(Platform):
                 ld_preload_str += pytorch_libgomp_so
                 os.environ["LD_PRELOAD"] = ld_preload_str
 
+        # To hint IPEX uses shared memory based AllReduce
         os.environ["LOCAL_WORLD_SIZE"] = str(
             vllm_config.parallel_config.tensor_parallel_size
         )

vllm/plugins/io_processors/interface.py

@@ -1,11 +1,12 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-import warnings
+
 from abc import ABC, abstractmethod
 from collections.abc import AsyncGenerator, Sequence
-from typing import Generic, TypeVar
+from typing import Any, Generic, TypeVar
 
 from vllm.config import VllmConfig
+from vllm.entrypoints.pooling.pooling.protocol import IOProcessorResponse
 from vllm.inputs.data import PromptType
 from vllm.outputs import PoolingRequestOutput
 from vllm.pooling_params import PoolingParams
@@ -17,68 +18,8 @@ IOProcessorOutput = TypeVar("IOProcessorOutput")
 
 class IOProcessor(ABC, Generic[IOProcessorInput, IOProcessorOutput]):
     def __init__(self, vllm_config: VllmConfig):
-        super().__init__()
-
         self.vllm_config = vllm_config
 
-    def parse_data(self, data: object) -> IOProcessorInput:
-        if callable(parse_request := getattr(self, "parse_request", None)):
-            warnings.warn(
-                "`parse_request` has been renamed to `parse_data`. "
-                "Please update your IO Processor Plugin to use the new name. "
-                "The old name will be removed in v0.19.",
-                DeprecationWarning,
-                stacklevel=2,
-            )
-
-            return parse_request(data)  # type: ignore
-
-        raise NotImplementedError
-
-    def merge_sampling_params(
-        self,
-        params: SamplingParams | None = None,
-    ) -> SamplingParams:
-        if callable(
-            validate_or_generate_params := getattr(
-                self, "validate_or_generate_params", None
-            )
-        ):
-            warnings.warn(
-                "`validate_or_generate_params` has been split into "
-                "`merge_sampling_params` and `merge_pooling_params`."
-                "Please update your IO Processor Plugin to use the new methods. "
-                "The old name will be removed in v0.19.",
-                DeprecationWarning,
-                stacklevel=2,
-            )
-
-            return validate_or_generate_params(params)  # type: ignore
-
-        return params or SamplingParams()
-
-    def merge_pooling_params(
-        self,
-        params: PoolingParams | None = None,
-    ) -> PoolingParams:
-        if callable(
-            validate_or_generate_params := getattr(
-                self, "validate_or_generate_params", None
-            )
-        ):
-            warnings.warn(
-                "`validate_or_generate_params` has been split into "
-                "`merge_sampling_params` and `merge_pooling_params`."
-                "Please update your IO Processor Plugin to use the new methods. "
-                "The old name will be removed in v0.19.",
-                DeprecationWarning,
-                stacklevel=2,
-            )
-
-            return validate_or_generate_params(params)  # type: ignore
-
-        return params or PoolingParams(task="plugin")
-
     @abstractmethod
     def pre_process(
         self,
@@ -118,4 +59,19 @@ class IOProcessor(ABC, Generic[IOProcessorInput, IOProcessorOutput]):
             [(i, item) async for i, item in model_output], key=lambda output: output[0]
         )
         collected_output = [output[1] for output in sorted_output]
-        return self.post_process(collected_output, request_id=request_id, **kwargs)
+        return self.post_process(collected_output, request_id, **kwargs)
+
+    @abstractmethod
+    def parse_request(self, request: Any) -> IOProcessorInput:
+        raise NotImplementedError
+
+    def validate_or_generate_params(
+        self, params: SamplingParams | PoolingParams | None = None
+    ) -> SamplingParams | PoolingParams:
+        return params or PoolingParams()
+
+    @abstractmethod
+    def output_to_response(
+        self, plugin_output: IOProcessorOutput
+    ) -> IOProcessorResponse:
+        raise NotImplementedError

vllm/renderers/__init__.py

@@ -1,8 +1,8 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 
-from .base import BaseRenderer
 from .params import ChatParams, TokenizeParams, merge_kwargs
+from .protocol import BaseRenderer
 from .registry import RendererRegistry, renderer_from_config
 
 __all__ = [

vllm/renderers/deepseek_v32.py

@@ -14,10 +14,10 @@ from vllm.tokenizers import cached_get_tokenizer
 from vllm.tokenizers.deepseek_v32 import DeepseekV32Tokenizer
 
 from ..tokenizers.hf import HfTokenizer
-from .base import BaseRenderer
 from .inputs import DictPrompt
 from .inputs.preprocess import parse_dec_only_prompt
 from .params import ChatParams
+from .protocol import BaseRenderer
 
 logger = init_logger(__name__)
 

vllm/renderers/grok2.py

@@ -13,10 +13,10 @@ from vllm.logger import init_logger
 from vllm.tokenizers import cached_get_tokenizer
 from vllm.tokenizers.grok2 import Grok2Tokenizer
 
-from .base import BaseRenderer
 from .inputs import DictPrompt
 from .inputs.preprocess import parse_dec_only_prompt
 from .params import ChatParams
+from .protocol import BaseRenderer
 
 logger = init_logger(__name__)
 

vllm/renderers/hf.py

@@ -32,10 +32,10 @@ from vllm.transformers_utils.chat_templates import get_chat_template_fallback_pa
 from vllm.transformers_utils.processor import cached_get_processor
 from vllm.utils.func_utils import supports_kw
 
-from .base import BaseRenderer
 from .inputs import DictPrompt
 from .inputs.preprocess import parse_dec_only_prompt
 from .params import ChatParams
+from .protocol import BaseRenderer
 
 if TYPE_CHECKING:
     from vllm.multimodal.inputs import MultiModalDataDict, MultiModalUUIDDict

vllm/renderers/mistral.py

@@ -15,10 +15,10 @@ from vllm.tokenizers import cached_get_tokenizer
 from vllm.tokenizers.mistral import MistralTokenizer
 from vllm.utils.async_utils import make_async
 
-from .base import BaseRenderer
 from .inputs import DictPrompt
 from .inputs.preprocess import parse_dec_only_prompt
 from .params import ChatParams
+from .protocol import BaseRenderer
 
 logger = init_logger(__name__)
 

vllm/renderers/registry.py

@@ -7,7 +7,7 @@ from vllm.logger import init_logger
 from vllm.tokenizers.registry import tokenizer_args_from_config
 from vllm.utils.import_utils import resolve_obj_by_qualname
 
-from .base import BaseRenderer
+from .protocol import BaseRenderer
 
 if TYPE_CHECKING:
     from vllm.config import ModelConfig

vllm/renderers/terratorch.py

@@ -12,10 +12,10 @@ from vllm.entrypoints.chat_utils import (
 from vllm.logger import init_logger
 from vllm.tokenizers import TokenizerLike
 
-from .base import BaseRenderer
 from .inputs import DictPrompt
 from .inputs.preprocess import parse_dec_only_prompt
 from .params import ChatParams
+from .protocol import BaseRenderer
 
 logger = init_logger(__name__)
 

vllm/transformers_utils/processors/__init__.py

@@ -10,7 +10,6 @@ reasons:
 
 from vllm.transformers_utils.processors.bagel import BagelProcessor
 from vllm.transformers_utils.processors.deepseek_vl2 import DeepseekVLV2Processor
-from vllm.transformers_utils.processors.funasr_processor import FunASRProcessor
 from vllm.transformers_utils.processors.hunyuan_vl import HunYuanVLProcessor
 from vllm.transformers_utils.processors.hunyuan_vl_image import HunYuanVLImageProcessor
 from vllm.transformers_utils.processors.ovis import OvisProcessor
@@ -19,7 +18,6 @@ from vllm.transformers_utils.processors.ovis2_5 import Ovis2_5Processor
 __all__ = [
     "BagelProcessor",
     "DeepseekVLV2Processor",
-    "FunASRProcessor",
     "HunYuanVLProcessor",
     "HunYuanVLImageProcessor",
     "OvisProcessor",

vllm/transformers_utils/processors/funasr_processor.py

@@ -1,504 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-
-
-import numpy as np
-import torch
-import torch.nn as nn
-import torchaudio.compliance.kaldi as kaldi
-from torch.nn.utils.rnn import pad_sequence
-from transformers import (
-    AutoFeatureExtractor,
-    AutoProcessor,
-    BatchFeature,
-)
-from transformers.feature_extraction_sequence_utils import SequenceFeatureExtractor
-from transformers.processing_utils import ProcessorMixin
-from transformers.utils import TensorType
-
-from vllm.logger import init_logger
-
-logger = init_logger(__name__)
-
-
-def apply_cmvn(inputs, cmvn):  # noqa
-    """
-    Apply CMVN with mvn data
-    """
-
-    device = inputs.device
-    # dtype = inputs.dtype
-    frame, dim = inputs.shape
-
-    means = cmvn[0:1, :dim]
-    vars = cmvn[1:2, :dim]
-    inputs += means.to(device)
-    inputs *= vars.to(device)
-
-    return inputs.type(torch.float32)
-
-
-def apply_lfr(inputs, lfr_m, lfr_n):
-    # LFR_inputs = []
-    T = inputs.shape[0]
-    T_lfr = int(np.ceil(T / lfr_n))
-    left_padding = inputs[0].repeat((lfr_m - 1) // 2, 1)
-    inputs = torch.vstack((left_padding, inputs))
-    T = T + (lfr_m - 1) // 2
-    feat_dim = inputs.shape[-1]
-    strides = (lfr_n * feat_dim, 1)
-    sizes = (T_lfr, lfr_m * feat_dim)
-    last_idx = (T - lfr_m) // lfr_n + 1
-    num_padding = lfr_m - (T - last_idx * lfr_n)
-    if num_padding > 0:
-        num_padding = (
-            (2 * lfr_m - 2 * T + (T_lfr - 1 + last_idx) * lfr_n)
-            / 2
-            * (T_lfr - last_idx)
-        )
-        inputs = torch.vstack([inputs] + [inputs[-1:]] * int(num_padding))
-    LFR_outputs = inputs.as_strided(sizes, strides)
-    return LFR_outputs.clone().type(torch.float32)
-
-
-def load_cmvn(cmvn_file):
-    with open(cmvn_file, encoding="utf-8") as f:
-        lines = f.readlines()
-    means_list = []
-    vars_list = []
-    for i in range(len(lines)):
-        line_item = lines[i].split()
-        if line_item[0] == "<AddShift>":
-            line_item = lines[i + 1].split()
-            if line_item[0] == "<LearnRateCoef>":
-                add_shift_line = line_item[3 : (len(line_item) - 1)]
-                means_list = list(add_shift_line)
-                continue
-        elif line_item[0] == "<Rescale>":
-            line_item = lines[i + 1].split()
-            if line_item[0] == "<LearnRateCoef>":
-                rescale_line = line_item[3 : (len(line_item) - 1)]
-                vars_list = list(rescale_line)
-                continue
-    means = np.array(means_list).astype(np.float32)
-    vars = np.array(vars_list).astype(np.float32)
-    cmvn = np.array([means, vars])
-    cmvn = torch.as_tensor(cmvn, dtype=torch.float32)
-    return cmvn
-
-
-class WavFrontend(nn.Module):
-    """Conventional frontend structure for ASR."""
-
-    def __init__(
-        self,
-        cmvn_file: str = "null",
-        fs: int = 16000,
-        window: str = "hamming",
-        n_mels: int = 80,
-        frame_length: int = 25,
-        frame_shift: int = 10,
-        filter_length_min: int = -1,
-        filter_length_max: int = -1,
-        lfr_m: int = 1,
-        lfr_n: int = 1,
-        dither: float = 1.0,
-        snip_edges: bool = True,
-        upsacle_samples: bool = True,
-        **kwargs,
-    ):
-        super().__init__()
-        self.fs = fs
-        self.window = window
-        self.n_mels = n_mels
-        self.frame_length = frame_length
-        self.frame_shift = frame_shift
-        self.filter_length_min = filter_length_min
-        self.filter_length_max = filter_length_max
-        self.lfr_m = lfr_m
-        self.lfr_n = lfr_n
-        self.cmvn_file = cmvn_file
-        self.dither = dither
-        self.snip_edges = snip_edges
-        self.upsacle_samples = upsacle_samples
-        self.cmvn = None if self.cmvn_file is None else load_cmvn(self.cmvn_file)
-
-    def output_size(self) -> int:
-        return self.n_mels * self.lfr_m
-
-    def forward(
-        self,
-        input: torch.Tensor,
-        input_lengths,
-        **kwargs,
-    ) -> tuple[torch.Tensor, torch.Tensor]:
-        batch_size = input.size(0)
-        feats = []
-        feats_lens = []
-        for i in range(batch_size):
-            waveform_length = input_lengths[i]
-            waveform = input[i][:waveform_length]
-            if self.upsacle_samples:
-                waveform = waveform * (1 << 15)
-            waveform = waveform.unsqueeze(0)
-            mat = kaldi.fbank(
-                waveform,
-                num_mel_bins=self.n_mels,
-                frame_length=min(self.frame_length, waveform_length / self.fs * 1000),
-                frame_shift=self.frame_shift,
-                dither=self.dither,
-                energy_floor=0.0,
-                window_type=self.window,
-                sample_frequency=self.fs,
-                snip_edges=self.snip_edges,
-            )
-
-            if self.lfr_m != 1 or self.lfr_n != 1:
-                mat = apply_lfr(mat, self.lfr_m, self.lfr_n)
-            if self.cmvn is not None:
-                mat = apply_cmvn(mat, self.cmvn)
-            feat_length = mat.size(0)
-            feats.append(mat)
-            feats_lens.append(feat_length)
-
-        feats_lens = torch.as_tensor(feats_lens)
-        if batch_size == 1:
-            feats_pad = feats[0][None, :, :]
-        else:
-            feats_pad = pad_sequence(feats, batch_first=True, padding_value=0.0)
-        return feats_pad, feats_lens
-
-    def forward_fbank(
-        self, input: torch.Tensor, input_lengths: torch.Tensor
-    ) -> tuple[torch.Tensor, torch.Tensor]:
-        batch_size = input.size(0)
-        feats = []
-        feats_lens = []
-        for i in range(batch_size):
-            waveform_length = input_lengths[i]
-            waveform = input[i][:waveform_length]
-            waveform = waveform * (1 << 15)
-            waveform = waveform.unsqueeze(0)
-            mat = kaldi.fbank(
-                waveform,
-                num_mel_bins=self.n_mels,
-                frame_length=self.frame_length,
-                frame_shift=self.frame_shift,
-                dither=self.dither,
-                energy_floor=0.0,
-                window_type=self.window,
-                sample_frequency=self.fs,
-            )
-
-            feat_length = mat.size(0)
-            feats.append(mat)
-            feats_lens.append(feat_length)
-
-        feats_lens = torch.as_tensor(feats_lens)
-        feats_pad = pad_sequence(feats, batch_first=True, padding_value=0.0)
-        return feats_pad, feats_lens
-
-    def forward_lfr_cmvn(
-        self, input: torch.Tensor, input_lengths: torch.Tensor
-    ) -> tuple[torch.Tensor, torch.Tensor]:
-        batch_size = input.size(0)
-        feats = []
-        feats_lens = []
-        for i in range(batch_size):
-            mat = input[i, : input_lengths[i], :]
-            if self.lfr_m != 1 or self.lfr_n != 1:
-                mat = apply_lfr(mat, self.lfr_m, self.lfr_n)
-            if self.cmvn is not None:
-                mat = apply_cmvn(mat, self.cmvn)
-            feat_length = mat.size(0)
-            feats.append(mat)
-            feats_lens.append(feat_length)
-
-        feats_lens = torch.as_tensor(feats_lens)
-        feats_pad = pad_sequence(feats, batch_first=True, padding_value=0.0)
-        return feats_pad, feats_lens
-
-
-class FunASRFeatureExtractor(SequenceFeatureExtractor):
-    r"""
-    Constructs a FunASR feature extractor.
-
-    This feature extractor inherits from [`~feature_extraction_sequence_
-        utils.SequenceFeatureExtractor`] which contains most of the main
-        methods. Users should refer to this superclass for more information
-        regarding those methods.
-
-    This class extracts mel-filter bank features from raw speech using a custom
-    numpy implementation of the `Short Time Fourier Transform` which should
-    match pytorch's `torch.stft` equivalent.
-
-    Args:
-        feature_size (`int`, *optional*, defaults to 80):
-            The feature dimension of the extracted features.
-        sampling_rate (`int`, *optional*, defaults to 16000):
-            The sampling rate at which the audio files should be digitalized
-            expressed in hertz (Hz).
-        hop_length (`int`, *optional*, defaults to 160):
-            Length of the overlapping windows for the STFT used to obtain the
-            Mel Frequency coefficients.
-        chunk_length (`int`, *optional*, defaults to 30):
-            The maximum number of chunks of `sampling_rate` samples used to
-            trim and pad longer or shorter audio sequences.
-        n_fft (`int`, *optional*, defaults to 400):
-            Size of the Fourier transform.
-        padding_value (`float`, *optional*, defaults to 0.0):
-            Padding value used to pad the audio. Should correspond to silences.
-        dither (`float`, *optional*, defaults to 0.0):
-            Adds dithering. In other words, adds a small Gaussian noise to each frame.
-            E.g. use 0.0001 to add dithering with a normal distribution centered
-            around 0.0 with standard deviation 0.0001 (assuming [-1,+1] range
-            of raw_speech). The value 0.0 means no dithering.
-            Dithering has similar effect as `spectrogram(mel_floor=...)`. It reduces
-            the high log_mel_fbank values for signals with hard-zero sections,
-            when VAD cutoff is present in the signal.
-    """
-
-    model_input_names = ["input_features"]
-
-    def __init__(
-        self,
-        feature_size=80,
-        sampling_rate=16000,
-        hop_length=160,
-        chunk_length=30,
-        n_fft=400,
-        padding_value=0.0,
-        dither=0.0,
-        return_attention_mask=False,
-        **kwargs,
-    ):
-        super().__init__(
-            feature_size=feature_size,
-            sampling_rate=sampling_rate,
-            padding_value=padding_value,
-            return_attention_mask=return_attention_mask,
-            **kwargs,
-        )
-        self.frontend_conf = kwargs.get("frontend_conf", {})
-        self.n_fft = n_fft
-        self.hop_length = hop_length
-        self.chunk_length = chunk_length
-        self.n_samples = chunk_length * sampling_rate
-        self.nb_max_frames = self.n_samples // hop_length
-        self.sampling_rate = sampling_rate
-        self.dither = dither
-
-    def extract_fbank(
-        self, data, data_len=None, data_type: str = "sound", frontend=None, **kwargs
-    ):
-        if isinstance(data, np.ndarray):
-            data = torch.from_numpy(data)
-            if len(data.shape) < 2:
-                data = data[None, :]  # data: [batch, N]
-            data_len = [data.shape[1]] if data_len is None else data_len
-        elif isinstance(data, torch.Tensor):
-            if len(data.shape) < 2:
-                data = data[None, :]  # data: [batch, N]
-            data_len = [data.shape[1]] if data_len is None else data_len
-        elif isinstance(data, (list, tuple)):
-            data_list, data_len = [], []
-            for data_i in data:
-                if isinstance(data_i, np.ndarray):
-                    data_i = torch.from_numpy(data_i)
-                data_list.append(data_i)
-                data_len.append(data_i.shape[0])
-            data = pad_sequence(data_list, batch_first=True)
-
-        data, data_len = frontend(data, data_len, **kwargs)
-
-        if isinstance(data_len, (list, tuple)):
-            data_len = torch.tensor([data_len])
-        return data.to(torch.float32), data_len.to(torch.int32)
-
-    def __call__(
-        self,
-        raw_speech: np.ndarray | list[float] | list[np.ndarray] | list[list[float]],
-        truncation: bool = True,
-        pad_to_multiple_of: int | None = None,
-        return_tensors: str | TensorType | None = None,
-        return_attention_mask: bool | None = None,
-        padding: str | None = "max_length",
-        max_length: int | None = None,
-        sampling_rate: int | None = None,
-        do_normalize: bool | None = None,
-        device: str | None = "cpu",
-        return_token_timestamps: bool | None = None,
-        **kwargs,
-    ) -> BatchFeature:
-        is_batched = isinstance(raw_speech, (list, tuple)) and (
-            isinstance(raw_speech[0], (np.ndarray, tuple, list))
-        )
-
-        if is_batched:
-            raw_speech = [
-                np.asarray([speech], dtype=np.float32).T for speech in raw_speech
-            ]
-        elif not is_batched and not isinstance(raw_speech, np.ndarray):
-            raw_speech = np.asarray(raw_speech, dtype=np.float32)
-        elif isinstance(raw_speech, np.ndarray) and raw_speech.dtype is np.dtype(
-            np.float64
-        ):
-            raw_speech = raw_speech.astype(np.float32)
-
-        if not is_batched:
-            raw_speech = [np.asarray([raw_speech]).T]
-
-        batched_speech = BatchFeature({"input_features": raw_speech})
-
-        padded_inputs = self.pad(
-            batched_speech,
-            padding=padding,
-            max_length=max_length if max_length else self.n_samples,
-            truncation=truncation,
-            pad_to_multiple_of=pad_to_multiple_of,
-            return_attention_mask=return_attention_mask or do_normalize,
-        )
-
-        input_features = padded_inputs.get("input_features").transpose(2, 0, 1)
-
-        self.frontend = WavFrontend(**self.frontend_conf)
-        input_features, speech_lengths = self.extract_fbank(
-            input_features[0],
-            data_type=kwargs.get("data_type", "sound"),
-            frontend=self.frontend,
-            is_final=True,
-        )
-        olens = 1 + (speech_lengths - 3 + 2 * 1) // 2
-        olens = 1 + (olens - 3 + 2 * 1) // 2
-        fake_token_len = (olens - 1) // 2 + 1
-        if isinstance(input_features[0], list):
-            padded_inputs["input_features"] = [
-                np.asarray(feature, dtype=np.float32) for feature in input_features
-            ]
-
-        else:
-            padded_inputs["input_features"] = input_features
-
-        if return_tensors is not None:
-            padded_inputs = padded_inputs.convert_to_tensors(return_tensors)
-
-        padded_inputs["speech_lengths"] = speech_lengths
-        padded_inputs["fake_token_len"] = fake_token_len
-
-        return padded_inputs
-
-
-class FunASRProcessor(ProcessorMixin):
-    r"""
-    Constructs a FunASR processor which wraps a FunASR feature extractor and
-    a FunASR tokenizer into a single processor.
-
-    [`FunASRProcessor`] offers all the functionalities of
-    [`FunASRFeatureExtractor`] and [`Qwen2Tokenizer`]. See the
-    [`~FunASRProcessor.__call__`] and [`~FunASRProcessor.decode`] for more
-    information.
-
-    Args:
-        feature_extractor (`FunASRFeatureExtractor`): An instance of
-            [`FunASRFeatureExtractor`].
-            The feature extractor is a required input.
-        tokenizer (`Qwen2Tokenizer`):
-            An instance of [`Qwen2Tokenizer`]. The tokenizer is a required
-            input.
-    """
-
-    feature_extractor_class = "FunASRFeatureExtractor"
-    tokenizer_class = ("Qwen2Tokenizer", "Qwen2TokenizerFast")
-
-    def __init__(
-        self,
-        feature_extractor,
-        tokenizer,
-        audio_token="<|AUDIO|>",
-    ):
-        super().__init__(feature_extractor, tokenizer)
-        self.current_processor = self.feature_extractor
-        self._in_target_context_manager = False
-        self.audio_token = (
-            tokenizer.audio_token if hasattr(tokenizer, "audio_token") else audio_token
-        )
-        self.audio_token_id = tokenizer.convert_tokens_to_ids(self.audio_token)
-
-    def get_decoder_prompt_ids(self, task=None, language=None, no_timestamps=True):
-        return self.tokenizer.get_decoder_prompt_ids(
-            task=task, language=language, no_timestamps=no_timestamps
-        )
-
-    def __call__(self, *args, **kwargs):
-        """
-        Forwards the `audio` argument to FunASRFeatureExtractor's
-        [`~FunASRFeatureExtractor.__call__`] and the `text` argument to
-        [`~Qwen2Tokenizer.__call__`]. Please refer to the docstring of the
-        above two methods for more information.
-        """
-        if self._in_target_context_manager:
-            return self.current_processor(*args, **kwargs)
-
-        audio = kwargs.pop("audio", None)
-        sampling_rate = kwargs.pop("sampling_rate", None)
-        text = kwargs.pop("text", None)
-        if len(args) > 0:
-            audio = args[0]
-            args = args[1:]
-
-        if text is None:
-            raise ValueError("You need to specify `text` input to process.")
-        elif isinstance(text, str):
-            text = [text]
-        elif not isinstance(text, list) and not isinstance(text[0], str):
-            raise ValueError(
-                "Invalid input text. Please provide a string, or a list of strings"
-            )
-
-        if audio is not None:
-            # ensure we have as much audios as audio tokens
-            num_audio_tokens = sum(sample.count(self.audio_token) for sample in text)
-            num_audios = 1 if type(audio) is np.ndarray else len(audio)
-            if num_audio_tokens != num_audios:
-                raise ValueError(
-                    f"Found {num_audio_tokens} {self.audio_token} token{'s' if num_audio_tokens > 1 else ''} in provided text but received {num_audios} audio{'s' if num_audios > 1 else ''}"  # noqa: E501
-                )
-            inputs = self.feature_extractor(
-                audio, *args, sampling_rate=sampling_rate, **kwargs
-            )
-
-            expanded_text = []
-            for sample in text:
-                replace_str = []
-                while self.audio_token in sample:
-                    num_audio_tokens = inputs["fake_token_len"].item()
-
-                    expanded_audio_token = self.audio_token * num_audio_tokens
-
-                    replace_str.append(expanded_audio_token)
-                    sample = sample.replace(self.audio_token, "<placeholder>", 1)
-
-                while "<placeholder>" in sample:
-                    sample = sample.replace("<placeholder>", replace_str.pop(0), 1)
-                expanded_text.append(sample)
-            text = expanded_text
-
-        if text is not None:
-            encodings = self.tokenizer(text, **kwargs)
-
-        if text is None:
-            return inputs
-
-        elif audio is None:
-            return encodings
-        else:
-            inputs["labels"] = encodings["input_ids"]
-
-            return inputs
-
-    def get_prompt_ids(self, text: str, return_tensors="np"):
-        return self.tokenizer.get_prompt_ids(text, return_tensors=return_tensors)
-
-
-AutoFeatureExtractor.register("FunASRFeatureExtractor", FunASRFeatureExtractor)
-AutoProcessor.register("FunASRProcessor", FunASRProcessor)

vllm/utils/collection_utils.py

@@ -51,6 +51,12 @@ def as_list(maybe_list: Iterable[T]) -> list[T]:
     return maybe_list if isinstance(maybe_list, list) else list(maybe_list)
 
 
+def as_iter(obj: T | Iterable[T]) -> Iterable[T]:
+    if isinstance(obj, str) or not isinstance(obj, Iterable):
+        return [obj]  # type: ignore[list-item]
+    return obj
+
+
 def is_list_of(
     value: object,
     typ: type[T] | tuple[type[T], ...],

vllm/v1/attention/backends/fa_utils.py

@@ -23,11 +23,12 @@ if current_platform.is_cuda():
 
 elif current_platform.is_xpu():
     from vllm import _custom_ops as ops
-    from vllm._xpu_ops import xpu_ops
 
     reshape_and_cache_flash = ops.reshape_and_cache_flash
-    flash_attn_varlen_func = xpu_ops.flash_attn_varlen_func  # type: ignore[assignment]
-    get_scheduler_metadata = xpu_ops.get_scheduler_metadata  # type: ignore[assignment]
+    from vllm._ipex_ops import ipex_ops
+
+    flash_attn_varlen_func = ipex_ops.flash_attn_varlen_func  # type: ignore[assignment]
+    get_scheduler_metadata = ipex_ops.get_scheduler_metadata  # type: ignore[assignment]
 elif current_platform.is_rocm():
     try:
         from flash_attn import flash_attn_varlen_func  # type: ignore[no-redef]
@@ -152,7 +153,7 @@ def is_flash_attn_varlen_func_available() -> bool:
 
     Platform-specific sources:
     - CUDA: vllm.vllm_flash_attn.flash_attn_varlen_func
-    - XPU: xpu_ops.flash_attn_varlen_func
+    - XPU: ipex_ops.flash_attn_varlen_func
     - ROCm: upstream flash_attn.flash_attn_varlen_func (if available)
 
     Note: This is separate from the AITER flash attention backend (rocm_aiter_fa.py)

vllm/v1/attention/backends/mla/indexer.py

@@ -268,39 +268,11 @@ class DeepseekV32IndexerMetadataBuilder(AttentionMetadataBuilder):
             num_reqs=reqs_end - reqs_start,
         )
 
-    def build_for_drafting(
-        self,
-        common_attn_metadata: CommonAttentionMetadata,
-        draft_index: int,
-    ) -> DeepseekV32IndexerMetadata:
-        """Build indexer metadata for draft model.
-
-        During drafting, each forward pass generates only 1 token per request,
-        so we set offsets=None (no speculative decoding offsets needed).
-        """
-        return self._build_impl(
-            common_attn_metadata=common_attn_metadata,
-            fast_build=True,
-            drafting=True,
-        )
-
     def build(
         self,
         common_prefix_len: int,
         common_attn_metadata: CommonAttentionMetadata,
         fast_build: bool = False,
-    ) -> DeepseekV32IndexerMetadata:
-        return self._build_impl(
-            common_attn_metadata=common_attn_metadata,
-            fast_build=fast_build,
-            drafting=False,
-        )
-
-    def _build_impl(
-        self,
-        common_attn_metadata: CommonAttentionMetadata,
-        fast_build: bool = False,
-        drafting: bool = False,
     ) -> DeepseekV32IndexerMetadata:
         num_reqs = common_attn_metadata.num_reqs
         num_tokens = common_attn_metadata.num_actual_tokens
@@ -359,18 +331,11 @@ class DeepseekV32IndexerMetadataBuilder(AttentionMetadataBuilder):
             # - top_k_per_row_decode wins for batch > 128 or seq_len <= 8K
             use_large_context_topk = batch_size <= 128 and _is_large_context
 
-            if drafting:
-                # During drafting, each forward generates 1 token per request,
-                # so no speculative offsets are needed.
-                offsets = None
+            next_n = 1 + self.num_speculative_tokens
+            if next_n > 1:
+                offsets = torch.arange(next_n, device=self.device, dtype=torch.int32)
             else:
-                next_n = 1 + self.num_speculative_tokens
-                if next_n > 1:
-                    offsets = torch.arange(
-                        next_n, device=self.device, dtype=torch.int32
-                    )
-                else:
-                    offsets = None
+                offsets = None
 
             seq_lens = common_attn_metadata.seq_lens[:num_decodes]
             if is_deep_gemm_supported():

vllm/v1/attention/ops/paged_attn.py

@@ -9,7 +9,7 @@ from vllm.platforms import current_platform
 if current_platform.is_cuda_alike():
     from vllm import _custom_ops as ops
 elif current_platform.is_xpu():
-    from vllm._xpu_ops import xpu_ops as ops  # type: ignore[no-redef]
+    from vllm._ipex_ops import ipex_ops as ops  # type: ignore[no-redef]
 
 
 class PagedAttention:

vllm/v1/engine/input_processor.py

@@ -72,15 +72,13 @@ class InputProcessor:
         self.mm_registry = mm_registry
         self.mm_processor_cache = mm_registry.processor_cache_from_config(vllm_config)
 
-        self.supports_mm_inputs = mm_registry.supports_multimodal_inputs(model_config)
-        self.mm_encoder_cache_size = 0
-        self.skip_prompt_length_check = False
-        if self.supports_mm_inputs:
+        self.mm_encoder_cache_size: int | None = None
+        if (
+            mm_registry.supports_multimodal_inputs(model_config)
+            and not model_config.skip_tokenizer_init
+        ):
             mm_budget = MultiModalBudget(vllm_config, mm_registry)
             self.mm_encoder_cache_size = mm_budget.encoder_cache_size
-            self.skip_prompt_length_check = (
-                mm_budget.processor.info.skip_prompt_length_check
-            )
             mm_budget.reset_cache()  # Not used anymore
 
         self.input_preprocessor = InputPreprocessor(
@@ -672,25 +670,76 @@ class InputProcessor:
             resumable=resumable,
         )
 
-    def _validate_prompt_len(
+    def _validate_model_inputs(
+        self, encoder_inputs: SingletonInputs | None, decoder_inputs: SingletonInputs
+    ):
+        if encoder_inputs is not None:
+            self._validate_model_input(encoder_inputs, prompt_type="encoder")
+
+        self._validate_model_input(decoder_inputs, prompt_type="decoder")
+
+    def _validate_model_input(
         self,
-        prompt_len: int,
+        prompt_inputs: SingletonInputs,
+        *,
         prompt_type: Literal["encoder", "decoder"],
     ):
-        if self.skip_prompt_length_check:
-            return
-
-        if prompt_len == 0 and prompt_type == "decoder":
-            raise ValueError(f"The {prompt_type} prompt cannot be empty")
-
         model_config = self.model_config
-        max_prompt_len = (
-            model_config.max_model_len
-            if prompt_type == "decoder"
-            else self.mm_encoder_cache_size
+
+        prompt_ids = (
+            None
+            if prompt_inputs["type"] == "embeds"
+            else prompt_inputs["prompt_token_ids"]
         )
+        prompt_embeds = (
+            prompt_inputs["prompt_embeds"]
+            if prompt_inputs["type"] == "embeds"
+            else None
+        )
+        prompt_len = length_from_prompt_token_ids_or_embeds(prompt_ids, prompt_embeds)
+        if not prompt_ids:
+            if prompt_type == "encoder" and model_config.is_multimodal_model:
+                pass  # Mllama may have empty encoder inputs for text-only data
+            elif prompt_inputs["type"] == "embeds":
+                pass  # Prompt embeds should not have prompt_ids.
+            else:
+                raise ValueError(f"The {prompt_type} prompt cannot be empty")
+
+        tokenizer = self.tokenizer
+        if tokenizer is not None:
+            max_input_id = max(prompt_ids or (), default=0)
+
+            # NOTE: tokenizer.max_token_id is the tokenizer’s vocab size while
+            # self.model_config.get_vocab_size() is the model’s vocab size.
+            # For Qwen3 models, the language model has extra tokens that do
+            # not exist in the tokenizer, and vice versa for multimodal
+            # placeholder tokens in some multimodal models.
+            # See https://github.com/QwenLM/Qwen3/issues/29#issuecomment-1933720399 # noqa: E501
+            # and https://github.com/vllm-project/vllm/pull/22471#discussion_r2312251421 # noqa: E501
+
+            # Here we take the max of the two to determine if a token id is
+            # truly out-of-vocabulary.
+            if max_input_id > max(
+                tokenizer.max_token_id, self.model_config.get_vocab_size() - 1
+            ):
+                raise ValueError(f"Token id {max_input_id} is out of vocabulary")
+
+        max_prompt_len = self.model_config.max_model_len
         if prompt_len > max_prompt_len:
-            if self.supports_mm_inputs:
+            if model_config.is_multimodal_model:
+                mm_registry = self.input_preprocessor.mm_registry
+                model_cls = mm_registry._get_model_cls(model_config)
+                factories = model_cls._processor_factory
+                ctx = mm_registry._create_processing_ctx(
+                    model_config,
+                    tokenizer=tokenizer,
+                )
+                mm_info = factories.info(ctx)
+
+                if mm_info.skip_prompt_length_check:
+                    return
+
+            if model_config.is_multimodal_model:
                 suggestion = (
                     "Make sure that `max_model_len` is no smaller than the "
                     "number of text tokens plus multimodal tokens. For image "
@@ -708,7 +757,17 @@ class InputProcessor:
                 f"longer than the maximum model length of {max_prompt_len}. "
                 f"{suggestion}"
             )
-        elif prompt_len == max_prompt_len and model_config.runner_type == "generate":
+
+            # TODO: Find out how many placeholder tokens are there so we can
+            # check that chunked prefill does not truncate them
+            # max_batch_len = self.scheduler_config.max_num_batched_tokens
+
+        if (
+            prompt_len == max_prompt_len
+            and prompt_type == "decoder"
+            and not model_config.is_multimodal_model
+            and self.model_config.runner_type != "pooling"
+        ):
             suggestion = (
                 "Make sure that `max_model_len` is no smaller than the "
                 "number of text tokens (prompt + requested output tokens)."
@@ -719,29 +778,11 @@ class InputProcessor:
                 f"model length of {max_prompt_len}. {suggestion}"
             )
 
-    def _validate_model_input(
-        self,
-        prompt_inputs: SingletonInputs,
-        prompt_type: Literal["encoder", "decoder"],
-    ) -> None:
-        model_config = self.model_config
-        tokenizer = self.tokenizer
-
-        prompt_ids = (
-            None
-            if prompt_inputs["type"] == "embeds"
-            else prompt_inputs["prompt_token_ids"]
-        )
-        prompt_embeds = (
-            prompt_inputs["prompt_embeds"]
-            if prompt_inputs["type"] == "embeds"
-            else None
-        )
-
-        prompt_len = length_from_prompt_token_ids_or_embeds(prompt_ids, prompt_embeds)
-        self._validate_prompt_len(prompt_len, prompt_type)
-
-        if prompt_inputs["type"] == "multimodal":
+        if (
+            prompt_type == "decoder"
+            and prompt_inputs["type"] == "multimodal"
+            and self.mm_encoder_cache_size is not None
+        ):
             decoder_mm_positions = prompt_inputs["mm_placeholders"]
             for modality, mm_positions in decoder_mm_positions.items():
                 for mm_position in mm_positions:
@@ -756,33 +797,6 @@ class InputProcessor:
                             f"by setting --limit-mm-per-prompt at startup."
                         )
 
-        if prompt_ids and tokenizer is not None:
-            max_input_id = max(prompt_ids, default=0)
-
-            # NOTE: tokenizer.max_token_id is the tokenizer’s vocab size while
-            # self.model_config.get_vocab_size() is the model’s vocab size.
-            # For Qwen3 models, the language model has extra tokens that do
-            # not exist in the tokenizer, and vice versa for multimodal
-            # placeholder tokens in some multimodal models.
-            # See https://github.com/QwenLM/Qwen3/issues/29#issuecomment-1933720399 # noqa: E501
-            # and https://github.com/vllm-project/vllm/pull/22471#discussion_r2312251421 # noqa: E501
-
-            # Here we take the max of the two to determine if a token id is
-            # truly out-of-vocabulary.
-            model_vocab_size = model_config.get_vocab_size()
-            if max_input_id > max(tokenizer.max_token_id, model_vocab_size - 1):
-                raise ValueError(f"Token id {max_input_id} is out of vocabulary")
-
-    def _validate_model_inputs(
-        self,
-        encoder_inputs: SingletonInputs | None,
-        decoder_inputs: SingletonInputs,
-    ):
-        if encoder_inputs is not None:
-            self._validate_model_input(encoder_inputs, prompt_type="encoder")
-
-        self._validate_model_input(decoder_inputs, prompt_type="decoder")
-
     def stat_mm_cache(self) -> MultiModalCacheStats | None:
         return self.input_preprocessor.stat_mm_cache()
 

vllm/v1/spec_decode/eagle.py

@@ -547,14 +547,10 @@ class SpecDecodeBaseProposer:
             num_tokens_unpadded=batch_size, num_tokens_padded=batch_size
         )
 
-        # Disable CUDA graphs for draft loop iterations.
-        # When CUDA graphs pad the batch (e.g., 100 → 104), the MLA attention
-        # backends call reshape_query_for_spec_decode(q, num_decodes) which
-        # fails when padded_size % num_decodes != 0. Even when divisible,
-        # padding positions contain invalid topk indices that cause CUDA
-        # illegal memory access in the FlashMLA kernel.
-        cudagraph_runtime_mode = CUDAGraphMode.NONE
-        input_batch_size = batch_size_dp_padded
+        cudagraph_runtime_mode, batch_desc = self.cudagraph_dispatcher.dispatch(
+            batch_size_dp_padded
+        )
+        input_batch_size = batch_desc.num_tokens
         if batch_size_across_dp is not None:
             batch_size_across_dp[self.dp_rank] = input_batch_size
 
@@ -657,17 +653,6 @@ class SpecDecodeBaseProposer:
             for layer_name in self.attn_layer_names:
                 per_layer_attn_metadata[layer_name] = attn_metadata
 
-            # Rebuild indexer/DSA attention metadata for models like GLM5
-            if self.draft_indexer_metadata_builder:
-                draft_indexer_metadata = (
-                    self.draft_indexer_metadata_builder.build_for_drafting(
-                        common_attn_metadata=common_attn_metadata,
-                        draft_index=token_index + 1,
-                    )
-                )
-                for layer_name in self.indexer_layer_names:
-                    per_layer_attn_metadata[layer_name] = draft_indexer_metadata
-
             # copy inputs to buffer for cudagraph
             self.input_ids[:batch_size] = input_ids
             self._set_positions(batch_size, clamped_positions)