add git

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

.buildkite/hardware_tests/amd.yaml

@@ -1,7 +1,6 @@
-group: Hardware - AMD Build 
+group: Hardware
 steps:
   - label: "AMD: :docker: build image"
-    key: image-build-amd
     depends_on: []
     device: amd_cpu
     no_plugin: true

.buildkite/test-amd.yaml

@@ -552,7 +552,7 @@ steps:
 - label: LoRA Test %N # 20min each
   timeout_in_minutes: 30
   mirror_hardwares: [amdexperimental]
-  agent_pool: mi325_1
+  agent_pool: mi325_8
   # grade: Blocking
   source_file_dependencies:
   - vllm/lora
@@ -648,7 +648,7 @@ steps:
 - label: Kernels Attention Test %N # 23min
   timeout_in_minutes: 35
   mirror_hardwares: [amdexperimental, amdproduction]
-  agent_pool: mi325_1
+  agent_pool: mi325_8
   # grade: Blocking
   source_file_dependencies:
   - csrc/attention/
@@ -663,7 +663,7 @@ steps:
 - label: Kernels Quantization Test %N # 64min
   timeout_in_minutes: 90
   mirror_hardwares: [amdexperimental]
-  agent_pool: mi325_1
+  agent_pool: mi325_8
   # grade: Blocking
   source_file_dependencies:
   - csrc/quantization/
@@ -676,7 +676,7 @@ steps:
 - label: Kernels MoE Test %N # 40min
   timeout_in_minutes: 60
   mirror_hardwares: [amdexperimental, amdproduction]
-  agent_pool: mi325_1
+  agent_pool: mi325_8
   # grade: Blocking
   source_file_dependencies:
   - csrc/quantization/cutlass_w8a8/moe/
@@ -839,7 +839,7 @@ steps:
 - label: Basic Models Tests (Extra Initialization) %N
   timeout_in_minutes: 45
   mirror_hardwares: [amdexperimental, amdproduction]
-  agent_pool: mi325_1
+  agent_pool: mi325_8
   # grade: Blocking
   torch_nightly: true
   source_file_dependencies:
@@ -901,7 +901,7 @@ steps:
 - label: Language Models Tests (Extra Standard) %N
   timeout_in_minutes: 45
   mirror_hardwares: [amdexperimental]
-  agent_pool: mi325_1
+  agent_pool: mi325_8
   # grade: Blocking
   torch_nightly: true
   source_file_dependencies:
@@ -922,7 +922,7 @@ steps:
 - label: Language Models Tests (Hybrid) %N
   timeout_in_minutes: 75
   mirror_hardwares: [amdexperimental]
-  agent_pool: mi325_1
+  agent_pool: mi325_8
   # grade: Blocking
   torch_nightly: true
   source_file_dependencies:

.buildkite/test_areas/basic_correctness.yaml

@@ -14,8 +14,3 @@ steps:
   - pytest -v -s basic_correctness/test_cumem.py
   - pytest -v -s basic_correctness/test_basic_correctness.py
   - pytest -v -s basic_correctness/test_cpu_offload.py
-  mirror:
-    amd:
-      device: mi325_1
-      depends_on:
-      - image-build-amd

.buildkite/test_areas/entrypoints.yaml

@@ -24,11 +24,6 @@ steps:
   - pytest -v -s entrypoints/llm --ignore=entrypoints/llm/test_generate.py --ignore=entrypoints/llm/test_collective_rpc.py
   - pytest -v -s entrypoints/llm/test_generate.py # it needs a clean process
   - pytest -v -s entrypoints/offline_mode # Needs to avoid interference with other tests
-  mirror:
-    amd:
-      device: mi325_1
-      depends_on:
-      - image-build-amd
 
 - label: Entrypoints Integration (API Server 1)
   timeout_in_minutes: 130

.buildkite/test_areas/models_basic.yaml

@@ -4,6 +4,7 @@ depends_on:
 steps:
 - label: Basic Models Tests (Initialization)
   timeout_in_minutes: 45
+  mirror_hardwares: [amdexperimental]
   torch_nightly: true
   source_file_dependencies:
   - vllm/
@@ -15,6 +16,7 @@ steps:
 
 - label: Basic Models Tests (Extra Initialization) %N
   timeout_in_minutes: 45
+  mirror_hardwares: [amdexperimental]
   torch_nightly: true
   source_file_dependencies:
   - vllm/model_executor/models/
@@ -36,12 +38,6 @@ steps:
   - tests/models/test_registry.py
   commands:
     - pytest -v -s models/test_terratorch.py models/test_transformers.py models/test_registry.py
-  mirror:
-    amd:
-      device: mi325_1
-      depends_on:
-      - image-build-amd
-    
 
 - label: Basic Models Test (Other CPU) # 5min
   depends_on: 

.buildkite/test_areas/models_language.yaml

@@ -4,6 +4,7 @@ depends_on:
 steps:
 - label: Language Models Tests (Standard)
   timeout_in_minutes: 25
+  mirror_hardwares: [amdexperimental]
   torch_nightly: true
   source_file_dependencies:
   - vllm/
@@ -15,6 +16,7 @@ steps:
 
 - label: Language Models Tests (Extra Standard) %N
   timeout_in_minutes: 45
+  mirror_hardwares: [amdexperimental]
   torch_nightly: true
   source_file_dependencies:
   - vllm/model_executor/models/
@@ -30,6 +32,7 @@ steps:
 
 - label: Language Models Tests (Hybrid) %N
   timeout_in_minutes: 75
+  mirror_hardwares: [amdexperimental]
   torch_nightly: true
   source_file_dependencies:
   - vllm/
@@ -45,6 +48,7 @@ steps:
 
 - label: Language Models Test (Extended Generation) # 80min
   timeout_in_minutes: 110
+  mirror_hardwares: [amdexperimental]
   optional: true
   source_file_dependencies:
   - vllm/
@@ -58,6 +62,7 @@ steps:
 
 - label: Language Models Test (PPL)
   timeout_in_minutes: 110
+  mirror_hardwares: [amdexperimental]
   optional: true
   source_file_dependencies:
   - vllm/
@@ -67,6 +72,7 @@ steps:
 
 - label: Language Models Test (Extended Pooling)  # 36min
   timeout_in_minutes: 50
+  mirror_hardwares: [amdexperimental]
   optional: true
   source_file_dependencies:
   - vllm/
@@ -76,6 +82,7 @@ steps:
 
 - label: Language Models Test (MTEB)
   timeout_in_minutes: 110
+  mirror_hardwares: [amdexperimental]
   optional: true
   source_file_dependencies:
   - vllm/

.buildkite/test_areas/samplers.yaml

@@ -12,10 +12,3 @@ steps:
   commands:
     - pytest -v -s samplers
     - VLLM_USE_FLASHINFER_SAMPLER=1 pytest -v -s samplers
-  mirror:
-    amd:
-      device: mi325_1
-      depends_on:
-      - image-build-amd
-      commands:
-      - pytest -v -s -m 'not skip_v1' samplers

.gitignore

@@ -238,6 +238,3 @@ ep_kernels_workspace/
 vllm/grpc/vllm_engine_pb2.py
 vllm/grpc/vllm_engine_pb2_grpc.py
 vllm/grpc/vllm_engine_pb2.pyi
-
-# Ignore generated cpu headers 
-csrc/cpu/cpu_attn_dispatch_generated.h

benchmarks/kernels/benchmark_cutlass_moe_fp8.py

@@ -11,7 +11,6 @@ import torch
 import vllm.model_executor.layers.fused_moe.modular_kernel as mk
 from tests.kernels.moe.utils import make_dummy_moe_config
 from vllm import _custom_ops as ops
-from vllm.model_executor.layers.fused_moe.activation import MoEActivation
 from vllm.model_executor.layers.fused_moe.config import fp8_w8a8_moe_quant_config
 from vllm.model_executor.layers.fused_moe.cutlass_moe import CutlassExpertsFp8
 from vllm.model_executor.layers.fused_moe.fused_moe import fused_experts, fused_topk
@@ -162,7 +161,7 @@ def bench_run(
                 w2_fp8q_cutlass,
                 topk_weights,
                 topk_ids,
-                activation=MoEActivation.SILU,
+                activation="silu",
                 global_num_experts=num_experts,
             )
     torch.cuda.synchronize()

benchmarks/kernels/benchmark_moe.py

@@ -16,7 +16,6 @@ import torch
 from ray.experimental.tqdm_ray import tqdm
 
 from vllm.model_executor.layers.fused_moe import fused_topk
-from vllm.model_executor.layers.fused_moe.activation import MoEActivation
 from vllm.model_executor.layers.fused_moe.config import (
     FusedMoEConfig,
     FusedMoEParallelConfig,
@@ -212,8 +211,7 @@ def benchmark_config(
                         hidden_dim=hidden_size,
                         intermediate_size_per_partition=shard_intermediate_size,
                         num_local_experts=num_experts,
-                        num_logical_experts=num_experts,
-                        activation=MoEActivation.SILU,
+                        activation="silu",
                         moe_parallel_config=FusedMoEParallelConfig.make_no_parallel(),
                         in_dtype=init_dtype,
                         routing_method=RoutingMethodType.TopK,
@@ -228,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,
@@ -239,7 +236,7 @@ def benchmark_config(
                 w2,
                 topk_weights,
                 topk_ids,
-                inplace=inplace,
+                inplace=True,
                 quant_config=quant_config,
             )
 

cmake/external_projects/vllm_flash_attn.cmake

@@ -38,7 +38,7 @@ else()
   FetchContent_Declare(
           vllm-flash-attn
           GIT_REPOSITORY https://github.com/vllm-project/flash-attention.git
-          GIT_TAG 5824e6e2008271063c3229ab3e7032bd74abbbc6
+          GIT_TAG 188be16520ceefdc625fdf71365585d2ee348fe2
           GIT_PROGRESS TRUE
           # Don't share the vllm-flash-attn build between build types
           BINARY_DIR ${CMAKE_BINARY_DIR}/vllm-flash-attn

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_fused_moe.cpp

@@ -147,7 +147,7 @@ void fused_moe_impl(scalar_t* __restrict__ output, scalar_t* __restrict__ input,
                     const int32_t token_num, const int32_t expert_num,
                     const int32_t topk_num, const int32_t input_size_13,
                     const int32_t output_size_13, const int32_t input_size_2,
-                    const int32_t output_size_2, const bool skip_weighted) {
+                    const int32_t output_size_2) {
   using scalar_vec_t = typename cpu_utils::VecTypeTrait<scalar_t>::vec_t;
   constexpr int32_t gemm_n_tile_size = gemm_t::NSize;
   constexpr int32_t gemm_m_tile_size = gemm_t::MaxMSize;
@@ -582,11 +582,6 @@ void fused_moe_impl(scalar_t* __restrict__ output, scalar_t* __restrict__ input,
         scalar_t* __restrict__ curr_output_buffer =
             output + token_id * output_size_2;
 
-        if (skip_weighted) {
-          // Only for topk_num == 1
-          *curr_weight = 1.0f;
-        }
-
         if (topk_num > 1) {
           {
             int32_t w2_output_idx = curr_expand_token_id_index_buffer[0];
@@ -704,7 +699,7 @@ void cpu_fused_moe(
     const std::optional<torch::Tensor>& w2_bias,  // [expert_num, output_size_2]
     const torch::Tensor& topk_weights,            // [token_num, k], float32
     const torch::Tensor& topk_id,                 // [token_num, k], int32
-    const bool skip_weighted, const std::string& act, const std::string& isa) {
+    const std::string& act, const std::string& isa) {
   const int32_t token_num = input.size(0);
   const int32_t input_size_13 = input.size(1);
   const int64_t input_stride = input.stride(0);
@@ -716,8 +711,6 @@ void cpu_fused_moe(
   const int32_t topk_num = topk_id.size(1);
   const FusedMOEAct act_type = get_act_type(act);
   cpu_utils::ISA isa_type = cpu_utils::get_isa(isa);
-  TORCH_CHECK(!skip_weighted || topk_num == 1,
-              "skip_weighted is only supported for topk=1 on CPU");
 
   VLLM_DISPATCH_FLOATING_TYPES(w13.scalar_type(), "cpu_fused_moe", [&]() {
     CPU_ISA_DISPATCH_IMPL(isa_type, [&]() {
@@ -728,7 +721,7 @@ void cpu_fused_moe(
           w2_bias.has_value() ? w2_bias->data_ptr<scalar_t>() : nullptr,
           topk_weights.data_ptr<float>(), topk_id.data_ptr<int32_t>(), act_type,
           token_num, expert_num, topk_num, input_size_13, output_size_13,
-          input_size_2, output_size_2, skip_weighted);
+          input_size_2, output_size_2);
     });
   });
 }

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;

csrc/cpu/torch_bindings.cpp

@@ -119,8 +119,8 @@ void cpu_fused_moe(torch::Tensor& output, const torch::Tensor& input,
                    const std::optional<torch::Tensor>& w13_bias,
                    const std::optional<torch::Tensor>& w2_bias,
                    const torch::Tensor& topk_weights,
-                   const torch::Tensor& topk_id, const bool skip_weighted,
-                   const std::string& act, const std::string& isa);
+                   const torch::Tensor& topk_id, const std::string& act,
+                   const std::string& isa);
 
 TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops) {
   // vLLM custom ops
@@ -320,7 +320,6 @@ TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops) {
   ops.def(
       "cpu_fused_moe(Tensor(a0!) output, Tensor input, Tensor w13, Tensor w2, "
       "Tensor? w13_bias, Tensor? w2_bias, Tensor topk_weights, Tensor topk_id, "
-      "bool skip_weighted, "
       "str act, str isa) -> ()");
   ops.impl("cpu_fused_moe", torch::kCPU, &cpu_fused_moe);
 #endif

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/features/quantization/README.md

@@ -48,7 +48,7 @@ th:not(:first-child) {
 |-----------------------|---------|----------|----------|-------|----------|-----------|-------------|-----------|
 | AWQ                   | ❌      | ✅︎       | ✅︎       | ✅︎    | ✅︎       | ❌         | ✅︎          | ✅︎        |
 | GPTQ                  | ✅︎      | ✅︎       | ✅︎       | ✅︎    | ✅︎       | ❌         | ✅︎          | ✅︎        |
-| Marlin (GPTQ/AWQ/FP8/FP4) | ❌      | ✅︎*       | ✅︎       | ✅︎    | ✅︎       | ❌         | ❌          | ❌        |
+| Marlin (GPTQ/AWQ/FP8) | ❌      | ❌       | ✅︎       | ✅︎    | ✅︎       | ❌         | ❌          | ❌        |
 | INT8 (W8A8)           | ❌      | ✅︎       | ✅︎       | ✅︎    | ✅︎       | ❌         | ❌          | ✅︎        |
 | FP8 (W8A8)            | ❌      | ❌       | ❌       | ✅︎    | ✅︎       | ✅︎         | ❌          | ❌        |
 | bitsandbytes          | ✅︎      | ✅︎       | ✅︎       | ✅︎    | ✅︎       | ❌         | ❌          | ❌        |
@@ -59,7 +59,6 @@ th:not(:first-child) {
 - ✅︎ indicates that the quantization method is supported on the specified hardware.
 - ❌ indicates that the quantization method is not supported on the specified hardware.
 - All Intel Gaudi quantization support has been migrated to [vLLM-Gaudi](https://github.com/vllm-project/vllm-gaudi).
-- *Turing does not support Marlin MXFP4.
 
 !!! note
     For information on quantization support on Google TPU, please refer to the [TPU-Inference Recommended Models and Features](https://docs.vllm.ai/projects/tpu/en/latest/recommended_models_features/) documentation.

docs/features/quantization/fp8.md

@@ -2,7 +2,7 @@
 
 vLLM supports FP8 (8-bit floating point) weight and activation quantization using hardware acceleration on GPUs such as Nvidia H100 and AMD MI300x.
 Currently, only Hopper and Ada Lovelace GPUs are officially supported for W8A8.
-Turing/Ampere GPUs are supported for W8A16 (weight-only FP8) utilizing Marlin kernels.
+Ampere GPUs are supported for W8A16 (weight-only FP8) utilizing Marlin kernels.
 Quantization of models with FP8 allows for a 2x reduction in model memory requirements and up to a 1.6x improvement in throughput with minimal impact on accuracy.
 
 Please visit the HF collection of [quantized FP8 checkpoints of popular LLMs ready to use with vLLM](https://huggingface.co/collections/neuralmagic/fp8-llms-for-vllm-666742ed2b78b7ac8df13127).
@@ -13,8 +13,8 @@ The FP8 types typically supported in hardware have two distinct representations,
 - **E5M2**: Consists of 1 sign bit, 5 exponent bits, and 2 bits of mantissa. It can store values up to +/-57344, +/- `inf`, and `nan`. The tradeoff for the increased dynamic range is lower precision of the stored values.
 
 !!! note
-    FP8 computation is supported on NVIDIA GPUs with compute capability >= 8.9 (Ada Lovelace, Hopper).
-    FP8 models will run on compute capability >= 7.5 (Turing) as weight-only W8A16, utilizing FP8 Marlin.
+    FP8 computation is supported on NVIDIA GPUs with compute capability > 8.9 (Ada Lovelace, Hopper).
+    FP8 models will run on compute capability > 8.0 (Ampere) as weight-only W8A16, utilizing FP8 Marlin.
 
 ## Installation
 

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)

mkdocs.yaml

@@ -63,9 +63,8 @@ plugins:
   - git-revision-date-localized:
       # exclude autogenerated files
       exclude:
-        - api/*
+        - argparse/*
         - examples/*
-        - generated/*
   - minify:
       minify_html: true
       minify_js: true
@@ -93,6 +92,7 @@ plugins:
               - "!.*_pb2_grpc"  # Exclude auto-generated gRPC stubs
             summary:
               modules: true
+            show_if_no_docstring: true
             show_signature_annotations: true
             separate_signature: true
             show_overloads: true

requirements/build.txt

@@ -9,5 +9,5 @@ wheel
 jinja2>=3.1.6
 regex
 build
-protobuf >= 5.29.6, !=6.30.*, !=6.31.*, !=6.32.*, !=6.33.0.*, !=6.33.1.*, !=6.33.2.*, !=6.33.3.*, !=6.33.4.*
+protobuf
 grpcio-tools==1.78.0 # Required for grpc entrypoints

requirements/common.txt

@@ -9,7 +9,7 @@ blake3
 py-cpuinfo
 transformers @ git+https://github.com/huggingface/transformers.git@main
 tokenizers >= 0.21.1  # Required for fast incremental detokenization.
-protobuf >= 5.29.6, !=6.30.*, !=6.31.*, !=6.32.*, !=6.33.0.*, !=6.33.1.*, !=6.33.2.*, !=6.33.3.*, !=6.33.4.* # Required by LlamaTokenizer, gRPC. CVE-2026-0994
+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.
 aiohttp >= 3.13.3
 openai >= 1.99.1  # For Responses API with reasoning content

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/compile/passes/test_fusion_attn.py

@@ -92,8 +92,6 @@ class AttentionQuantPatternModel(torch.nn.Module):
     def build_attn_metadata(self, batch_size: int) -> AttentionMetadata:
         """Initialize attention metadata."""
 
-        # TODO (Rohan138) reuse utils from vllm/v1/worker/gpu/attn_utils.py
-
         # Create common attn metadata
         batch_spec = BatchSpec(seq_lens=[1] * batch_size, query_lens=[1] * batch_size)
         common_attn_metadata = create_common_attn_metadata(
@@ -102,31 +100,58 @@ class AttentionQuantPatternModel(torch.nn.Module):
 
         max_blocks = (max(batch_spec.seq_lens) + self.block_size - 1) // self.block_size
         num_blocks = batch_size * max_blocks
+        backend = self.attn.backend
 
-        # Fetch the attention backend and kv cache shape and stride order
-        attn_backend = self.attn.attn_backend
-        kv_cache_shape = attn_backend.get_kv_cache_shape(
-            num_blocks, self.block_size, self.num_kv_heads, self.head_size
-        )
-        try:
-            kv_cache_stride_order = attn_backend.get_kv_cache_stride_order()
-        except (AttributeError, NotImplementedError):
-            kv_cache_stride_order = tuple(range(len(kv_cache_shape)))
-
-        kv_cache_shape = tuple(kv_cache_shape[i] for i in kv_cache_stride_order)
-        inv_order = [
-            kv_cache_stride_order.index(i) for i in range(len(kv_cache_stride_order))
-        ]
-
-        # Create dummy KV cache
-        raw_tensor = torch.zeros(
-            2 * num_blocks * self.block_size * self.num_kv_heads * self.head_size,
-            dtype=self.kv_cache_dtype,
-            device=self.device,
-        )
-        raw_tensor = raw_tensor.view(kv_cache_shape)
-        kv_cache = raw_tensor.permute(*inv_order)
-
+        # TODO(luka) use get_kv_cache_stride_order
+        # Create dummy KV cache for the selected backend
+        if backend == AttentionBackendEnum.ROCM_ATTN:
+            # k/v as 1st dimention
+            # HND: [num_blocks, num_kv_heads, block_size, head_size]
+            kv_cache = torch.zeros(
+                2,
+                num_blocks,
+                self.num_kv_heads,
+                self.block_size,
+                self.head_size,
+                dtype=self.kv_cache_dtype,
+                device=self.device,
+            )
+        elif backend == AttentionBackendEnum.ROCM_AITER_UNIFIED_ATTN:
+            # k/v as 1st dimention
+            # NHD: [num_blocks, block_size, num_kv_heads, head_size]
+            kv_cache = torch.zeros(
+                2,
+                num_blocks,
+                self.block_size,
+                self.num_kv_heads,
+                self.head_size,
+                dtype=self.kv_cache_dtype,
+                device=self.device,
+            )
+        elif backend == AttentionBackendEnum.TRITON_ATTN:
+            # k/v as 2nd dimention
+            # NHD: [num_blocks, block_size, num_kv_heads, head_size]
+            kv_cache = torch.zeros(
+                num_blocks,
+                2,
+                self.num_kv_heads,
+                self.block_size,
+                self.head_size,
+                dtype=self.kv_cache_dtype,
+                device=self.device,
+            )
+        elif backend == AttentionBackendEnum.FLASHINFER:
+            kv_cache = torch.zeros(
+                num_blocks,
+                2,
+                self.num_kv_heads,
+                self.block_size,
+                self.head_size,
+                dtype=self.kv_cache_dtype,
+                device=self.device,
+            ).permute(0, 1, 3, 2, 4)
+        else:
+            raise ValueError(f"Unsupported backend: {backend}")
         self.attn.kv_cache = [kv_cache]
 
         # Build attn metadata

tests/config/test_config_generation.py

@@ -80,19 +80,12 @@ def test_ray_runtime_env(monkeypatch: pytest.MonkeyPatch):
     ray.shutdown()
 
 
-def test_unrecognized_env(monkeypatch):
+def test_unrecognized_env():
     import os
 
-    from vllm.envs import environment_variables
-
-    # Remove any existing unrecognized VLLM env vars that might interfere
-    for env in list(os.environ):
-        if env.startswith("VLLM_") and env not in environment_variables:
-            monkeypatch.delenv(env, raising=False)
-
     # Test that if fail_on_environ_validation is True, then an error
     # is raised when an unrecognized vLLM environment variable is set
-    monkeypatch.setenv("VLLM_UNRECOGNIZED_ENV_VAR", "some_value")
+    os.environ["VLLM_UNRECOGNIZED_ENV_VAR"] = "some_value"
     engine_args = EngineArgs(
         fail_on_environ_validation=True,
     )
@@ -104,7 +97,7 @@ def test_unrecognized_env(monkeypatch):
     engine_args.create_engine_config()
 
     # Test that when the unrecognized env var is removed, no error is raised
-    monkeypatch.delenv("VLLM_UNRECOGNIZED_ENV_VAR")
+    os.environ.pop("VLLM_UNRECOGNIZED_ENV_VAR", None)
     engine_args = EngineArgs(
         fail_on_environ_validation=True,
     )

tests/entrypoints/openai/responses/test_harmony.py

@@ -1302,17 +1302,16 @@ async def test_system_prompt_override(client: OpenAI, model_name: str):
         # Message structure may vary, skip this specific check
         pass
 
-    custom_system_prompt_2 = (
-        "You are a helpful assistant that always responds in exactly 5 words."
-    )
-
     # Test 3: Test with different custom system prompt
     response_2 = await client.responses.create(
         model=model_name,
         input=[
             {
                 "role": "system",
-                "content": custom_system_prompt_2,
+                "content": (
+                    "You are a helpful assistant that always "
+                    "responds in exactly 5 words."
+                ),
             },
             {"role": "user", "content": "What is the weather like?"},
         ],
@@ -1329,27 +1328,3 @@ async def test_system_prompt_override(client: OpenAI, model_name: str):
     assert 3 <= word_count <= 8, (
         f"Expected around 5 words, got {word_count} words: {response_2.output_text}"
     )
-
-    # Test 4: Test with structured content
-    response_3 = await client.responses.create(
-        model=model_name,
-        input=[
-            {
-                "role": "system",
-                "content": [{"type": "input_text", "text": custom_system_prompt_2}],
-            },
-            {"role": "user", "content": "What is the weather like?"},
-        ],
-        temperature=0.0,
-    )
-
-    assert response_3 is not None
-    assert response_3.status == "completed"
-    assert response_3.output_text is not None
-
-    # Count words in response (approximately, allowing for punctuation)
-    word_count = len(response_3.output_text.split())
-    # Allow some flexibility (4-7 words) since the model might not be perfectly precise
-    assert 3 <= word_count <= 8, (
-        f"Expected around 5 words, got {word_count} words: {response_3.output_text}"
-    )

tests/entrypoints/openai/test_realtime_validation.py

@@ -129,5 +129,5 @@ async def test_multi_chunk_streaming(
                 " First words I spoke in the original phonograph."
                 " A little piece of practical poetry. Mary had a little lamb,"
                 " it sleeps with quite a flow, and everywhere that Mary went,"
-                " the lamb was sure to go."
+                " the lamb was sure to go"
             )

tests/evals/gsm8k/configs/moe-refactor/Llama-4-Scout-BF16-fi-cutlass.yaml

@@ -5,5 +5,3 @@ num_fewshot: 5
 server_args: "--enforce-eager --max-model-len 8192 --tensor-parallel-size 2 --enable-expert-parallel"
 env:
   VLLM_USE_FLASHINFER_MOE_FP16: "1"
-  VLLM_FLASHINFER_MOE_BACKEND: "throughput"
-

tests/evals/gsm8k/configs/moe-refactor/Mixtral-8x7B-BF16-fi-cutlass.yaml

@@ -5,4 +5,3 @@ num_fewshot: 5
 server_args: "--enforce-eager --max-model-len 8192 --tensor-parallel-size 2 --enable-expert-parallel"
 env:
   VLLM_USE_FLASHINFER_MOE_FP16: "1"
-  VLLM_FLASHINFER_MOE_BACKEND: "throughput"

tests/kernels/moe/modular_kernel_tools/common.py

@@ -22,7 +22,6 @@ from vllm.distributed import (
 )
 from vllm.forward_context import set_forward_context
 from vllm.model_executor.layers.fused_moe import fused_topk
-from vllm.model_executor.layers.fused_moe.activation import MoEActivation
 from vllm.model_executor.layers.fused_moe.all2all_utils import (
     maybe_make_prepare_finalize,
 )
@@ -600,7 +599,7 @@ def make_modular_kernel(
         moe_parallel_config=moe_parallel_config,
         in_dtype=config.dtype,
         max_num_tokens=next_power_of_2(config.M),
-        activation=MoEActivation.SILU,
+        activation="silu",
         device=vllm_config.device_config.device,
         routing_method=RoutingMethodType.DeepSeekV3,
     )

tests/kernels/moe/test_cpu_fused_moe.py

@@ -6,7 +6,6 @@ import torch
 
 from tests.kernels.allclose_default import get_default_atol, get_default_rtol
 from vllm._custom_ops import cpu_fused_moe, cpu_prepack_moe_weight
-from vllm.model_executor.layers.fused_moe.activation import MoEActivation
 from vllm.model_executor.layers.fused_moe.cpu_fused_moe import _CPU_MOE_ACT_FN
 from vllm.platforms import current_platform
 from vllm.utils.torch_utils import set_random_seed
@@ -20,7 +19,7 @@ EXPERT_NUM = [
 HIDDEN_DIM = [128, 2880]
 INTERMEDIATE_DIM = [128, 2880]
 BATCH_SIZE = [1, 64, 256]
-ACT = [MoEActivation.SILU, MoEActivation.SWIGLUOAI]
+ACT = ["silu", "swigluoai"]
 USE_BIAS = [True, False]
 ISA = ["amx", "vec"] if torch._C._cpu._is_amx_tile_supported() else ["vec"]
 DTYPE = [torch.bfloat16]
@@ -34,7 +33,7 @@ def ref_fused_moe(
     w2_bias: torch.Tensor | None,
     topk_weights: torch.Tensor,
     topk_ids: torch.Tensor,
-    activation: MoEActivation,
+    activation: str,
 ) -> torch.Tensor:
     len_experts = w13.size(0)
 
@@ -104,7 +103,7 @@ def test_cpu_fused_moe(
     intermediate_size: int,
     use_bias: bool,
     dtype: torch.dtype,
-    act: MoEActivation,
+    act: str,
     isa: str,
 ):
     set_random_seed(0)
@@ -154,7 +153,7 @@ def test_cpu_fused_moe(
         w2_bias,
         topk_weight,
         topk_ids,
-        act.value,
+        act,
         isa,
     )
 

tests/kernels/moe/test_cutlass_moe.py

@@ -12,7 +12,6 @@ from tests.kernels.moe.utils import make_dummy_moe_config
 from vllm import _custom_ops as ops
 from vllm.config import ParallelConfig, VllmConfig, set_current_vllm_config
 from vllm.model_executor.layers.fused_moe import fused_experts, fused_topk
-from vllm.model_executor.layers.fused_moe.activation import MoEActivation
 from vllm.model_executor.layers.fused_moe.config import (
     FUSED_MOE_UNQUANTIZED_CONFIG,
     FusedMoEQuantConfig,
@@ -532,7 +531,7 @@ def test_run_cutlass_moe_fp8(
         c_strides1 = torch.full((e,), 2 * n, device="cuda", dtype=torch.int64)
         c_strides2 = torch.full((e,), k, device="cuda", dtype=torch.int64)
 
-        activation = MoEActivation.SILU
+        activation = "silu"
         a1q, a1q_scale = moe_kernel_quantize_input(
             mt.a, mt.a_scale, torch.float8_e4m3fn, per_act_token
         )

tests/kernels/moe/test_deepep_deepgemm_moe.py

@@ -16,7 +16,6 @@ from typing_extensions import ParamSpec
 
 from vllm.config import VllmConfig, set_current_vllm_config
 from vllm.forward_context import set_forward_context
-from vllm.model_executor.layers.fused_moe.activation import MoEActivation
 from vllm.model_executor.layers.fused_moe.config import (
     FusedMoEQuantConfig,
     fp8_w8a8_moe_quant_config,
@@ -325,7 +324,7 @@ def deepep_deepgemm_moe_impl(
             w2=w2,
             topk_weights=test_tensors.topk_weights,
             topk_ids=test_tensors.topk,
-            activation=MoEActivation.SILU,
+            activation="silu",
             global_num_experts=num_experts,
             expert_map=build_expert_map(),
             apply_router_weight_on_input=False,

tests/kernels/moe/test_deepep_moe.py

@@ -15,7 +15,6 @@ from vllm import _custom_ops as ops
 from vllm.config import VllmConfig, set_current_vllm_config
 from vllm.model_executor.layers.activation import SiluAndMul
 from vllm.model_executor.layers.fused_moe import TritonExperts
-from vllm.model_executor.layers.fused_moe.activation import MoEActivation
 from vllm.model_executor.layers.fused_moe.config import (
     FusedMoEQuantConfig,
 )
@@ -261,7 +260,7 @@ def deep_ep_moe_impl(
             w2=w2,
             topk_weights=topk_weights_chunk,
             topk_ids=topk_chunk,
-            activation=MoEActivation.SILU,
+            activation="silu",
             global_num_experts=num_experts,
             expert_map=build_expert_map(),
             apply_router_weight_on_input=False,

tests/kernels/moe/test_flashinfer.py

@@ -7,7 +7,6 @@ import torch
 
 import vllm.model_executor.layers.fused_moe.modular_kernel as mk
 from vllm.config import ParallelConfig, VllmConfig, set_current_vllm_config
-from vllm.model_executor.layers.fused_moe.activation import MoEActivation
 from vllm.model_executor.layers.fused_moe.config import (
     FusedMoEConfig,
     FusedMoEParallelConfig,
@@ -94,14 +93,9 @@ class TestData:
 
     @staticmethod
     def make_moe_tensors_8bit(
-        m: int,
-        k: int,
-        n: int,
-        e: int,
-        is_trtllm: bool,
-        activation: MoEActivation = MoEActivation.SILU,
+        m: int, k: int, n: int, e: int, is_trtllm: bool, activation: str = "silu"
     ) -> "TestData":
-        is_gated = activation.is_gated
+        is_gated = activation != "relu2_no_mul"
 
         hidden_states = torch.randn((m, k), device="cuda", dtype=torch.bfloat16) / 10
         w13 = torch.randn(
@@ -200,7 +194,7 @@ def test_flashinfer_per_tensor_moe_fp8_no_graph(
             topk_weights=topk_weights,
             topk_ids=topk_ids,
             inplace=False,
-            activation=MoEActivation.SILU,
+            activation="silu",
             global_num_experts=e,
             expert_map=None,
             apply_router_weight_on_input=True,
@@ -225,19 +219,21 @@ def test_flashinfer_per_tensor_moe_fp8_no_graph(
 @pytest.mark.parametrize("m,n,k", MNK_FACTORS)
 @pytest.mark.parametrize("e", NUM_EXPERTS)
 @pytest.mark.parametrize("topk", TOP_KS)
-@pytest.mark.parametrize("activation", [MoEActivation.SILU, MoEActivation.RELU2_NO_MUL])
+@pytest.mark.parametrize("activation", ["silu", "relu2_no_mul"])
 def test_flashinfer_cutlass_moe_fp8_no_graph(
     m: int,
     n: int,
     k: int,
     e: int,
     topk: int,
-    activation: MoEActivation,
+    activation: str,
     monkeypatch,
     workspace_init,
 ):
     set_random_seed(7)
     monkeypatch.setenv("VLLM_FUSED_MOE_CHUNK_SIZE", "8192")
+    assert activation in ["silu", "relu2_no_mul"]
+    is_act_and_mul = activation == "silu_and_mul"
     with set_current_vllm_config(vllm_config):
         td = TestData.make_moe_tensors_8bit(
             m, k, n, e, is_trtllm=False, activation=activation
@@ -291,12 +287,11 @@ def test_flashinfer_cutlass_moe_fp8_no_graph(
             hidden_dim=k,
             intermediate_size_per_partition=n,
             num_local_experts=e,
-            num_logical_experts=e,
             activation=activation,
             device="cuda",
             moe_parallel_config=FusedMoEParallelConfig.make_no_parallel(),
             in_dtype=torch.bfloat16,
-            is_act_and_mul=activation.is_gated,
+            is_act_and_mul=is_act_and_mul,
             routing_method=RoutingMethodType.TopK,
         )
 
@@ -323,44 +318,3 @@ def test_flashinfer_cutlass_moe_fp8_no_graph(
         torch.testing.assert_close(
             output, flashinfer_cutlass_output, atol=5.5e-2, rtol=1e-2
         )
-
-
-@pytest.mark.parametrize(
-    "num_experts,intermediate,hidden",
-    [
-        (8, 2048, 1536),
-        (64, 4096, 4096),
-    ],
-)
-def test_convert_moe_weights_to_flashinfer_trtllm_block_layout(
-    num_experts, intermediate, hidden
-):
-    from vllm.model_executor.layers.quantization.utils.flashinfer_utils import (
-        convert_moe_weights_to_flashinfer_trtllm_block_layout,
-    )
-
-    w13 = torch.randn(
-        (num_experts, 2 * intermediate, hidden), dtype=torch.bfloat16, device="cuda"
-    )
-    w2 = torch.randn(
-        (num_experts, hidden, intermediate), dtype=torch.bfloat16, device="cuda"
-    )
-
-    cache: dict[torch.Size, torch.Tensor] = {}
-    w13_converted, w2_converted = convert_moe_weights_to_flashinfer_trtllm_block_layout(
-        cache, w13, w2
-    )
-
-    assert w13_converted.ndim == 4, (
-        f"Expected 4D tensor, got shape {w13_converted.shape}"
-    )
-    assert w2_converted.ndim == 4, f"Expected 4D tensor, got shape {w2_converted.shape}"
-
-    assert w13_converted.numel() == w13.numel(), "W13 element count should be preserved"
-    assert w2_converted.numel() == w2.numel(), "W2 element count should be preserved"
-
-    assert w13_converted.dtype == torch.bfloat16
-    assert w2_converted.dtype == torch.bfloat16
-
-    assert w13_converted.shape[0] == num_experts
-    assert w2_converted.shape[0] == num_experts

tests/kernels/moe/test_flashinfer_moe.py

@@ -13,7 +13,6 @@ from tests.kernels.utils import torch_moe
 from vllm import _custom_ops as ops
 from vllm.config import ParallelConfig, VllmConfig, set_current_vllm_config
 from vllm.model_executor.layers.fused_moe import fused_topk
-from vllm.model_executor.layers.fused_moe.activation import MoEActivation
 from vllm.model_executor.layers.fused_moe.config import (
     FusedMoEConfig,
     FusedMoEParallelConfig,
@@ -55,7 +54,7 @@ MNK_FACTORS = [
 @pytest.mark.parametrize("e", [40, 64, 256])
 @pytest.mark.parametrize("topk", [1, 6, 8])
 @pytest.mark.parametrize("dtype", [torch.bfloat16])
-@pytest.mark.parametrize("activation", [MoEActivation.SILU, MoEActivation.RELU2_NO_MUL])
+@pytest.mark.parametrize("activation", ["silu_and_mul", "relu2"])
 @torch.inference_mode()
 def test_flashinfer_fp4_moe_no_graph(
     m: int,
@@ -64,7 +63,7 @@ def test_flashinfer_fp4_moe_no_graph(
     e: int,
     topk: int,
     dtype: torch.dtype,
-    activation: MoEActivation,
+    activation: str,
     workspace_init,
 ):
     set_random_seed(7)
@@ -74,7 +73,7 @@ def test_flashinfer_fp4_moe_no_graph(
         a = torch.randn((m, k), device="cuda", dtype=dtype) / 10
 
         quant_blocksize = 16
-        is_gated_act = activation.is_gated
+        is_gated_act = activation == "silu_and_mul"
 
         w1_q, w2_q, quant_config = make_test_quant_config(
             e,
@@ -98,7 +97,6 @@ def test_flashinfer_fp4_moe_no_graph(
             hidden_dim=k,
             intermediate_size_per_partition=n,
             num_local_experts=e,
-            num_logical_experts=e,
             activation=activation,
             device="cuda",
             moe_parallel_config=FusedMoEParallelConfig.make_no_parallel(),
@@ -113,13 +111,15 @@ def test_flashinfer_fp4_moe_no_graph(
             inplace=False,
         )
 
+        fi_activation = {"silu_and_mul": "silu", "relu2": "relu2_no_mul"}[activation]
+
         flashinfer_output = flashinfer_experts(
             hidden_states=a,
             w1=w1_q,
             w2=w2_q,
             topk_weights=topk_weights,
             topk_ids=topk_ids,
-            activation=activation,
+            activation=fi_activation,
         )
 
         # Reference check:

tests/kernels/moe/test_modular_oai_triton_moe.py

@@ -7,7 +7,6 @@ Test modular OAI Triton MoE
 import pytest
 import torch
 
-from vllm.model_executor.layers.fused_moe.activation import MoEActivation
 from vllm.utils.import_utils import has_triton_kernels
 
 if not has_triton_kernels():
@@ -193,7 +192,7 @@ def oai_triton_moe_impl(
         w2=w2,
         topk_weights=topk_weights,
         topk_ids=topk_ids,
-        activation=MoEActivation.SWIGLUOAI,
+        activation="swigluoai",
         global_num_experts=num_experts,
         expert_map=None,
         apply_router_weight_on_input=False,

tests/kernels/moe/test_moe.py

@@ -29,7 +29,6 @@ from vllm.config import VllmConfig, set_current_vllm_config
 from vllm.distributed.parallel_state import init_distributed_environment
 from vllm.forward_context import get_forward_context, set_forward_context
 from vllm.model_executor.layers.fused_moe import (
-    MoEActivation,
     fused_topk,
 )
 from vllm.model_executor.layers.fused_moe.config import (
@@ -1156,10 +1155,7 @@ def test_fused_marlin_moe_with_bias(m):
 @pytest.mark.parametrize("m", [1, 64, 256])
 @pytest.mark.parametrize("n,k", [(1024, 1024), (2048, 2048)])
 @pytest.mark.parametrize("e,topk", [(8, 2), (64, 4)])
-@pytest.mark.parametrize("activation", [MoEActivation.RELU2_NO_MUL])
-def test_fused_marlin_moe_non_gated(
-    m: int, n: int, k: int, e: int, topk: int, activation: MoEActivation
-):
+def test_fused_marlin_moe_non_gated(m: int, n: int, k: int, e: int, topk: int):
     """Test Marlin MoE with non-gated activation (relu2_no_mul).
 
     Non-gated activations like relu2 don't have the gate-up projection pattern,
@@ -1202,7 +1198,7 @@ def test_fused_marlin_moe_non_gated(
             w2_data.w_ref,
             score,
             topk,
-            activation=activation,
+            activation="relu2",
         )
 
     marlin_output = fused_marlin_moe(
@@ -1227,7 +1223,7 @@ def test_fused_marlin_moe_non_gated(
         w2_zeros=w2_data.zeros,
         quant_type_id=quant_type.id,
         is_k_full=is_k_full,
-        activation=activation,
+        activation="relu2_no_mul",
     )
 
     torch.testing.assert_close(marlin_output, torch_output, atol=1e-1, rtol=0)
@@ -1334,18 +1330,9 @@ def test_moe_sum(m: int, topk: int, k: int, dtype: torch.dtype):
 @pytest.mark.parametrize("topk", [2])
 @pytest.mark.parametrize("dtype", [torch.float32, torch.bfloat16])
 @pytest.mark.parametrize("with_bias", [False, True])
-@pytest.mark.parametrize("activation", [MoEActivation.SILU])
+@pytest.mark.parametrize("activation", ["silu"])
 @pytest.mark.skipif(not current_platform.is_cpu(), reason="CPU only test")
-def test_cpu_fused_moe_basic(
-    m: int,
-    n: int,
-    k: int,
-    e: int,
-    topk: int,
-    dtype: torch.dtype,
-    with_bias: bool,
-    activation: MoEActivation,
-):
+def test_cpu_fused_moe_basic(m, n, k, e, topk, dtype, with_bias, activation):
     from vllm.model_executor.layers.fused_moe.cpu_fused_moe import CPUFusedMOE
 
     device = "cpu"
@@ -1571,104 +1558,3 @@ def test_batched_fused_marlin_moe(
     marlin_output = br.run(a, kwargs)
 
     torch.testing.assert_close(marlin_output, ref_marlin_output, atol=1e-3, rtol=0)
-
-
-@pytest.mark.parametrize("m,n,k", [(32, 1024, 1024)])
-@pytest.mark.parametrize("e,topk", [(8, 2)])
-@pytest.mark.parametrize("dtype", [torch.bfloat16])
-@pytest.mark.skipif(
-    not current_platform.is_device_capability_family(100),
-    reason="TRTLLM backend test only runs on Blackwell GPUs (SM10x).",
-)
-def test_unquantized_bf16_flashinfer_trtllm_backend(
-    m: int,
-    n: int,
-    k: int,
-    e: int,
-    topk: int,
-    dtype: torch.dtype,
-    monkeypatch,
-    workspace_init,
-):
-    """
-    Test BF16 unquantized MoE with FlashInfer TRTLLM backend.
-    """
-    set_random_seed(7)
-
-    monkeypatch.setenv("VLLM_USE_FLASHINFER_MOE_FP16", "1")
-
-    from vllm.model_executor.layers.fused_moe.config import (
-        FusedMoEConfig,
-        FusedMoEParallelConfig,
-        RoutingMethodType,
-    )
-    from vllm.model_executor.layers.fused_moe.oracle.unquantized import (
-        UnquantizedMoeBackend,
-    )
-    from vllm.model_executor.layers.fused_moe.unquantized_fused_moe_method import (
-        UnquantizedFusedMoEMethod,
-    )
-
-    # Setup test data
-    a = torch.randn((m, k), device="cuda", dtype=dtype) / 10
-    w1 = torch.randn((e, 2 * n, k), device="cuda", dtype=dtype) / 10
-    w2 = torch.randn((e, k, n), device="cuda", dtype=dtype) / 10
-    router_logits = torch.randn((m, e), device="cuda", dtype=dtype)
-
-    moe_config = FusedMoEConfig(
-        num_experts=e,
-        experts_per_token=topk,
-        hidden_dim=k,
-        intermediate_size_per_partition=n,
-        num_local_experts=e,
-        num_logical_experts=e,
-        activation="silu",
-        device="cuda",
-        moe_parallel_config=FusedMoEParallelConfig.make_no_parallel(),
-        in_dtype=dtype,
-        is_act_and_mul=True,
-        routing_method=RoutingMethodType.Renormalize,
-        max_num_tokens=m,
-    )
-
-    with set_current_vllm_config(vllm_config):
-        quant_method = UnquantizedFusedMoEMethod(moe_config)
-
-        # Verify TRTLLM backend was selected
-        assert (
-            quant_method.unquantized_backend == UnquantizedMoeBackend.FLASHINFER_TRTLLM
-        ), f"Expected FLASHINFER_TRTLLM backend, got {quant_method.unquantized_backend}"
-
-        # Verify it's using monolithic path
-        assert quant_method.is_monolithic, (
-            "FLASHINFER_TRTLLM backend should use monolithic forward"
-        )
-        layer = torch.nn.Module()
-        layer.w13_weight = Parameter(w1.clone(), requires_grad=False)
-        layer.w2_weight = Parameter(w2.clone(), requires_grad=False)
-        layer.global_num_experts = e
-        layer.local_num_experts = e
-        layer.top_k = topk
-        layer.num_expert_group = 1
-        layer.topk_group = 1
-        layer.intermediate_size_per_partition = n
-        layer.ep_rank = 0
-        layer.activation = "silu"
-        layer.e_score_correction_bias = None
-        layer.routing_method_type = RoutingMethodType.Renormalize
-
-        quant_method.process_weights_after_loading(layer)
-
-        trtllm_output = quant_method.forward_monolithic_cuda(
-            layer=layer,
-            x=a,
-            router_logits=router_logits,
-        )
-
-        # Compute torch baseline
-        w1_original = w1.clone()
-        w2_original = w2.clone()
-        baseline_output = torch_moe(a, w1_original, w2_original, router_logits, topk)
-
-    close = torch.isclose(trtllm_output, baseline_output, atol=1e-1, rtol=0.85)
-    assert close.float().mean() > 0.925

tests/kernels/moe/test_pplx_cutlass_moe.py

@@ -9,7 +9,6 @@ from tests.kernels.utils import torch_experts
 from vllm import _custom_ops as ops
 from vllm.config import VllmConfig, set_current_vllm_config
 from vllm.model_executor.layers.fused_moe import fused_topk
-from vllm.model_executor.layers.fused_moe.activation import MoEActivation
 from vllm.model_executor.layers.fused_moe.config import (
     FusedMoEConfig,
     FusedMoEParallelConfig,
@@ -148,9 +147,8 @@ def pplx_cutlass_moe(
             hidden_dim=hidden_dim,
             intermediate_size_per_partition=intermediate_dim,
             num_local_experts=num_local_experts,
-            num_logical_experts=num_experts,
             moe_parallel_config=FusedMoEParallelConfig.make_no_parallel(),
-            activation=MoEActivation.SILU,
+            activation="silu",
             in_dtype=torch.bfloat16,
             device="cuda",
             routing_method=RoutingMethodType.Llama4,

tests/kernels/moe/test_triton_moe_no_act_mul.py

@@ -11,11 +11,15 @@ import pytest
 import torch
 
 from tests.kernels.moe.utils import make_dummy_moe_config
-from vllm.model_executor.layers.fused_moe.activation import MoEActivation
 from vllm.model_executor.layers.fused_moe.config import (
     FUSED_MOE_UNQUANTIZED_CONFIG,
 )
 from vllm.model_executor.layers.fused_moe.fused_moe import TritonExperts
+from vllm.model_executor.layers.fused_moe.utils import (
+    GELU_NO_MUL,
+    RELU2_NO_MUL,
+    SILU_NO_MUL,
+)
 from vllm.platforms import current_platform
 
 # Test parameters
@@ -24,11 +28,7 @@ N_SIZES = [128, 256]
 K_SIZES = [64, 128]
 TOPK_VALUES = [1, 2]
 NUM_EXPERTS = 8
-NO_MUL_ACTIVATIONS = [
-    MoEActivation.SILU_NO_MUL,
-    MoEActivation.GELU_NO_MUL,
-    MoEActivation.RELU2_NO_MUL,
-]
+NO_MUL_ACTIVATIONS = [SILU_NO_MUL, GELU_NO_MUL, RELU2_NO_MUL]
 
 
 def make_test_tensors(
@@ -73,7 +73,7 @@ def test_triton_experts_no_mul_activation(
     n: int,
     k: int,
     topk: int,
-    activation: MoEActivation,
+    activation: str,
 ):
     hidden_states, w1, w2, topk_weights, topk_ids = make_test_tensors(
         m, n, k, NUM_EXPERTS, topk
@@ -161,11 +161,11 @@ def test_workspace_shapes_no_mul_vs_gated():
     )
 
     ws1_no_mul, _, out_no_mul = experts.workspace_shapes(
-        M, N, K, topk, 8, 8, None, MoEActivation.SILU_NO_MUL
+        M, N, K, topk, 8, 8, None, SILU_NO_MUL
     )
 
     ws1_gated, _, out_gated = experts.workspace_shapes(
-        M, N, K, topk, 8, 8, None, MoEActivation.SILU
+        M, N, K, topk, 8, 8, None, "silu"
     )
 
     # For no_mul: activation_out_dim = N
@@ -202,10 +202,10 @@ def test_adjust_n_for_activation():
     N = 256
 
     # Gated activations should return N // 2
-    assert experts.adjust_N_for_activation(N, MoEActivation.SILU) == N // 2
-    assert experts.adjust_N_for_activation(N, MoEActivation.GELU) == N // 2
+    assert experts.adjust_N_for_activation(N, "silu") == N // 2
+    assert experts.adjust_N_for_activation(N, "gelu") == N // 2
 
     # Non-gated activations should return N
-    assert experts.adjust_N_for_activation(N, MoEActivation.SILU_NO_MUL) == N
-    assert experts.adjust_N_for_activation(N, MoEActivation.GELU_NO_MUL) == N
-    assert experts.adjust_N_for_activation(N, MoEActivation.RELU2_NO_MUL) == N
+    assert experts.adjust_N_for_activation(N, SILU_NO_MUL) == N
+    assert experts.adjust_N_for_activation(N, GELU_NO_MUL) == N
+    assert experts.adjust_N_for_activation(N, RELU2_NO_MUL) == N

tests/kernels/moe/test_unquantized_backend_selection.py

@@ -1,132 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-from unittest.mock import patch
-
-import pytest
-
-from tests.kernels.moe.utils import make_dummy_moe_config
-from vllm.model_executor.layers.fused_moe.oracle.unquantized import (
-    UnquantizedMoeBackend,
-    select_unquantized_moe_backend,
-)
-
-
-@pytest.mark.parametrize(
-    "platform_method,expected_backend",
-    [
-        ("is_cuda", UnquantizedMoeBackend.TRITON),  # Default CUDA without FlashInfer
-        ("is_rocm", UnquantizedMoeBackend.TRITON),
-        ("is_cpu", UnquantizedMoeBackend.CPU),
-        ("is_xpu", UnquantizedMoeBackend.XPU),
-        ("is_tpu", UnquantizedMoeBackend.TPU),
-        ("is_out_of_tree", UnquantizedMoeBackend.OOT),
-    ],
-)
-@patch(
-    "vllm.model_executor.layers.fused_moe.oracle.unquantized.has_flashinfer",
-    return_value=False,
-)
-def test_select_default_backend_by_platform(
-    mock_has_flashinfer,
-    monkeypatch,
-    platform_method,
-    expected_backend,
-):
-    """Test backend selection for different platforms."""
-    with patch(
-        "vllm.model_executor.layers.fused_moe.oracle.unquantized.current_platform"
-    ) as mock_platform:
-        # Set all platform checks to False
-        mock_platform.is_cuda.return_value = False
-        mock_platform.is_rocm.return_value = False
-        mock_platform.is_cpu.return_value = False
-        mock_platform.is_xpu.return_value = False
-        mock_platform.is_tpu.return_value = False
-        mock_platform.is_out_of_tree.return_value = False
-
-        # Set only the specified platform to True
-        getattr(mock_platform, platform_method).return_value = True
-
-        moe_config = make_dummy_moe_config()
-        selected_backend = select_unquantized_moe_backend(
-            moe_config=moe_config,
-            use_ep=False,
-            use_dp=False,
-        )
-
-        assert selected_backend == expected_backend
-
-
-@patch(
-    "vllm.model_executor.layers.fused_moe.oracle.unquantized.has_flashinfer",
-    return_value=True,
-)
-@patch(
-    "vllm.model_executor.layers.fused_moe.oracle.unquantized.is_supported_config_trtllm_bf16",
-    return_value=(True, None),
-)
-def test_select_cuda_flashinfer_trtllm_backend(
-    mock_has_flashinfer, mock_is_supported_trtllm, monkeypatch
-):
-    """Test CUDA backend selection when FlashInfer TRTLLM is available and enabled."""
-    with patch(
-        "vllm.model_executor.layers.fused_moe.oracle.unquantized.current_platform"
-    ) as mock_platform:
-        # Set as CUDA platform
-        mock_platform.is_cuda.return_value = True
-        mock_platform.is_rocm.return_value = False
-        mock_platform.is_cpu.return_value = False
-        mock_platform.is_xpu.return_value = False
-        mock_platform.is_tpu.return_value = False
-        mock_platform.is_out_of_tree.return_value = False
-
-        monkeypatch.setenv("VLLM_USE_FLASHINFER_MOE_FP16", "1")
-
-        moe_config = make_dummy_moe_config()
-
-        selected_backend = select_unquantized_moe_backend(
-            moe_config=moe_config,
-            use_ep=True,
-            use_dp=False,
-        )
-
-        assert selected_backend == UnquantizedMoeBackend.FLASHINFER_TRTLLM
-
-
-@patch(
-    "vllm.model_executor.layers.fused_moe.oracle.unquantized.has_flashinfer",
-    return_value=True,
-)
-@patch(
-    "vllm.model_executor.layers.fused_moe.oracle.unquantized.is_supported_config_trtllm_bf16",
-    return_value=(False, None),
-)
-def test_select_cuda_flashinfer_cutlass_backend(
-    mock_has_flashinfer, mock_is_supported_trtllm, monkeypatch
-):
-    """Test CUDA backend selection when FlashInfer TRTLLM is not available
-    and FlashInfer CUTLASS is available."""
-    with patch(
-        "vllm.model_executor.layers.fused_moe.oracle.unquantized.current_platform"
-    ) as mock_platform:
-        # Set as CUDA platform with Hopper capability
-        mock_platform.is_cuda.return_value = True
-        mock_platform.is_rocm.return_value = False
-        mock_platform.is_cpu.return_value = False
-        mock_platform.is_xpu.return_value = False
-        mock_platform.is_tpu.return_value = False
-        mock_platform.is_out_of_tree.return_value = False
-        mock_platform.has_device_capability.return_value = True  # SM90+
-
-        # Enable FlashInfer via env var
-        monkeypatch.setenv("VLLM_USE_FLASHINFER_MOE_FP16", "1")
-
-        moe_config = make_dummy_moe_config()
-
-        selected_backend = select_unquantized_moe_backend(
-            moe_config=moe_config,
-            use_ep=True,  # CUTLASS requires EP
-            use_dp=False,  # CUTLASS doesn't support DP
-        )
-
-        assert selected_backend == UnquantizedMoeBackend.FLASHINFER_CUTLASS

tests/kernels/moe/utils.py

@@ -12,7 +12,6 @@ from vllm.model_executor.layers.fused_moe import (
     fused_experts,
     fused_topk,
 )
-from vllm.model_executor.layers.fused_moe.activation import MoEActivation
 from vllm.model_executor.layers.fused_moe.config import (
     FusedMoEConfig,
     FusedMoEParallelConfig,
@@ -55,7 +54,7 @@ def make_dummy_moe_config(
         num_local_experts=num_experts,
         num_logical_experts=num_experts,
         moe_parallel_config=FusedMoEParallelConfig.make_no_parallel(),
-        activation=MoEActivation.SILU,
+        activation="silu",
         in_dtype=in_dtype,
         device="cuda",
         routing_method=RoutingMethodType.TopK,

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/kernels/utils.py

@@ -15,7 +15,6 @@ from torch._prims_common import TensorLikeType
 from tests.kernels.quant_utils import native_w8a8_block_matmul
 from vllm.model_executor.custom_op import op_registry
 from vllm.model_executor.layers.activation import SiluAndMul
-from vllm.model_executor.layers.fused_moe.activation import MoEActivation
 from vllm.model_executor.layers.fused_moe.utils import moe_kernel_quantize_input
 from vllm.utils.torch_utils import make_tensor_with_pad
 from vllm.v1.attention.backend import AttentionType
@@ -841,7 +840,7 @@ def torch_experts(
     per_act_token_quant=False,
     block_shape: list[int] | None = None,
     apply_router_weights_on_input: bool = False,
-    activation: MoEActivation = MoEActivation.SILU,
+    activation: str = "silu_and_mul",
 ) -> torch.Tensor:
     assert (
         global_num_experts == -1
@@ -884,7 +883,7 @@ def torch_experts(
 
     f32 = torch.float32
 
-    act = op_registry[activation.custom_op_name]
+    act = op_registry[activation]
 
     for i in range(num_experts):
         mask = topk_ids == i
@@ -974,7 +973,7 @@ def torch_moe(
     b_bias2: torch.Tensor | None = None,
     global_num_experts: int = -1,
     expert_map: torch.Tensor | None = None,
-    activation: MoEActivation = MoEActivation.SILU,
+    activation: str = "silu_and_mul",
 ) -> torch.Tensor:
     score = torch.softmax(score, dim=-1, dtype=torch.float32)
     topk_weight, topk_ids = torch.topk(score, topk)

tests/models/multimodal/processing/test_common.py

@@ -108,7 +108,6 @@ _ADD_SPECIAL_TOKENS_OVERRIDES = {
     "paligemma": False,
     "ultravox": False,
     "whisper": False,
-    "lfm2_vl": False,
 }
 
 _IGNORE_MM_KEYS = {

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
     ),
@@ -725,6 +721,7 @@ _MULTIMODAL_EXAMPLE_MODELS = {
     "Gemma3nForConditionalGeneration": _HfExamplesInfo("google/gemma-3n-E2B-it"),
     "GlmAsrForConditionalGeneration": _HfExamplesInfo(
         "zai-org/GLM-ASR-Nano-2512",
+        trust_remote_code=True,
         min_transformers_version="5.0.0",
     ),
     "GraniteVision": _HfExamplesInfo("ibm-granite/granite-vision-3.3-2b"),

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_blackwell_moe.py

@@ -178,11 +178,3 @@ def test_gptoss_eager(monkeypatch: pytest.MonkeyPatch):
         hf_overrides=HF_OVERRIDE_TEXT,
         extra_args=["--enforce-eager"],
     )
-
-
-## Qwen3 Next ##
-
-
-def test_qwen3_next_bf16_moe_flashinfer_trtllm(monkeypatch: pytest.MonkeyPatch):
-    monkeypatch.setenv("VLLM_USE_FLASHINFER_MOE_FP16", "1")
-    can_initialize("Qwen/Qwen3-Next-80B-A3B-Instruct", hf_overrides=HF_OVERRIDE_TEXT)

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/cudagraph/test_cudagraph_dispatch.py

@@ -1,6 +1,5 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-from dataclasses import replace
 from unittest.mock import MagicMock, patch
 
 import pytest
@@ -133,39 +132,36 @@ class TestCudagraphDispatcher:
 
         # Test dispatch logic
         # 1. non-uniform batch, size in cudagraph size list
-        # FULL mode uses exact keys with num_reqs set
-        desc_full_with_reqs = BatchDescriptor(num_tokens=8, num_reqs=8, uniform=False)
-        # PIECEWISE mode uses relaxed keys with num_reqs=None
-        desc_piecewise = BatchDescriptor(num_tokens=8, num_reqs=None, uniform=False)
+        desc_full_exact = BatchDescriptor(
+            num_tokens=8,
+            uniform=False,
+        )
         rt_mode, key = dispatcher.dispatch(
             num_tokens=8, uniform_decode=False, has_lora=False
         )
         if cudagraph_mode_str == "FULL":
             assert rt_mode == CUDAGraphMode.FULL
-            assert key == desc_full_with_reqs
+            assert key == desc_full_exact
         elif cudagraph_mode_str in ["FULL_AND_PIECEWISE", "PIECEWISE"]:
             assert rt_mode == CUDAGraphMode.PIECEWISE
-            assert key == desc_piecewise
+            assert key == desc_full_exact
         else:
             assert rt_mode == CUDAGraphMode.NONE
 
         # 2. uniform decode batch, size in cudagraph size list
         desc_uniform_exact = BatchDescriptor(num_tokens=8, num_reqs=8, uniform=True)
-        desc_non_uniform = BatchDescriptor(num_tokens=8, num_reqs=8, uniform=False)
         rt_mode, key = dispatcher.dispatch(
             num_tokens=8, uniform_decode=True, has_lora=False
         )
         if cudagraph_mode_str == "FULL":
-            # Pure FULL mode uses non-uniform keys for all batches
             assert rt_mode == CUDAGraphMode.FULL
-            assert key == desc_non_uniform
+            assert key == desc_uniform_exact.relax_for_mixed_batch_cudagraphs()
         elif cudagraph_mode_str in ["FULL_DECODE_ONLY", "FULL_AND_PIECEWISE"]:
-            # These modes have separate uniform decode keys
             assert rt_mode == CUDAGraphMode.FULL
             assert key == desc_uniform_exact
         elif cudagraph_mode_str == "PIECEWISE":
             assert rt_mode == CUDAGraphMode.PIECEWISE
-            assert key == replace(desc_uniform_exact, num_reqs=None, uniform=False)
+            assert key == desc_uniform_exact.relax_for_mixed_batch_cudagraphs()
         else:
             assert rt_mode == CUDAGraphMode.NONE
 
@@ -184,7 +180,7 @@ class TestCudagraphDispatcher:
 
         if "PIECEWISE" in cudagraph_mode_str:  # string contains check
             assert rt_mode == CUDAGraphMode.PIECEWISE
-            assert key == replace(desc_full_exact, num_reqs=None, uniform=False)
+            assert key == desc_full_exact.relax_for_mixed_batch_cudagraphs()
         else:
             assert rt_mode == CUDAGraphMode.NONE
 

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,
     )

tools/pre_commit/shellcheck.baseline

@@ -1,89 +0,0 @@
-benchmarks/auto_tune/auto_tune.sh:SC2034
-benchmarks/auto_tune/auto_tune.sh:SC2086
-benchmarks/auto_tune/batch_auto_tune.sh:SC2086
-benchmarks/run_structured_output_benchmark.sh:SC2028
-benchmarks/run_structured_output_benchmark.sh:SC2034
-benchmarks/run_structured_output_benchmark.sh:SC2086
-.buildkite/image_build/image_build_cpu_arm64.sh:SC2086
-.buildkite/image_build/image_build_cpu.sh:SC2086
-.buildkite/image_build/image_build_hpu.sh:SC2086
-.buildkite/lm-eval-harness/run-lm-eval-chartqa-vllm-vlm-baseline.sh:SC2086
-.buildkite/lm-eval-harness/run-lm-eval-mmlupro-vllm-baseline.sh:SC2034
-.buildkite/performance-benchmarks/scripts/run-performance-benchmarks.sh:SC2027
-.buildkite/performance-benchmarks/scripts/run-performance-benchmarks.sh:SC2086
-.buildkite/performance-benchmarks/scripts/run-performance-benchmarks.sh:SC2126
-.buildkite/scripts/annotate-rocm-release.sh:SC2086
-.buildkite/scripts/cache-rocm-base-wheels.sh:SC2012
-.buildkite/scripts/cherry-pick-from-milestone.sh:SC2064
-.buildkite/scripts/hardware_ci/run-cpu-test-ppc64le.sh:SC2086
-.buildkite/scripts/hardware_ci/run-cpu-test.sh:SC2086
-.buildkite/scripts/hardware_ci/run-hpu-test.sh:SC2086
-.buildkite/scripts/hardware_ci/run-npu-test.sh:SC1090
-.buildkite/scripts/hardware_ci/run-npu-test.sh:SC2006
-.buildkite/scripts/hardware_ci/run-npu-test.sh:SC2086
-.buildkite/scripts/hardware_ci/run-npu-test.sh:SC2181
-.buildkite/scripts/hardware_ci/run-xpu-test.sh:SC2086
-.buildkite/scripts/push-nightly-builds.sh:SC2086
-.buildkite/scripts/run-multi-node-test.sh:SC2086
-.buildkite/scripts/run-multi-node-test.sh:SC2089
-.buildkite/scripts/run-multi-node-test.sh:SC2090
-.buildkite/scripts/run-prime-rl-test.sh:SC2086
-.buildkite/scripts/scheduled_integration_test/deepseek_v2_lite_ep_eplb.sh:SC2086
-.buildkite/scripts/scheduled_integration_test/qwen30b_a3b_fp8_block_ep_eplb.sh:SC2086
-.buildkite/scripts/scheduled_integration_test/qwen3_next_mtp_async_eplb.sh:SC2086
-.buildkite/scripts/tpu/docker_run_bm.sh:SC1090
-.buildkite/scripts/tpu/docker_run_bm.sh:SC2086
-.buildkite/scripts/tpu/run_bm.sh:SC2034
-.buildkite/scripts/tpu/run_bm.sh:SC2086
-.buildkite/scripts/upload-nightly-wheels.sh:SC2086
-.buildkite/scripts/upload-nightly-wheels.sh:SC2115
-.buildkite/scripts/upload-nightly-wheels.sh:SC2236
-.buildkite/scripts/upload-release-wheels-pypi.sh:SC2086
-.buildkite/scripts/upload-rocm-wheels.sh:SC2012
-examples/online_serving/disaggregated_encoder/disagg_1e1p1d_example.sh:SC2086
-examples/online_serving/disaggregated_encoder/disagg_1e1pd_example.sh:SC2086
-examples/online_serving/disaggregated_prefill.sh:SC2086
-examples/online_serving/disaggregated_serving/kv_events.sh:SC2086
-examples/online_serving/disaggregated_serving/mooncake_connector/run_mooncake_connector.sh:SC2046
-examples/online_serving/disaggregated_serving/mooncake_connector/run_mooncake_connector.sh:SC2086
-examples/online_serving/disaggregated_serving/mooncake_connector/run_mooncake_connector.sh:SC2317
-examples/online_serving/disaggregated_serving_p2p_nccl_xpyd/disagg_example_p2p_nccl_xpyd.sh:SC2046
-examples/online_serving/disaggregated_serving_p2p_nccl_xpyd/disagg_example_p2p_nccl_xpyd.sh:SC2086
-examples/online_serving/disaggregated_serving_p2p_nccl_xpyd/disagg_example_p2p_nccl_xpyd.sh:SC2317
-examples/online_serving/elastic_ep/bench.sh:SC2086
-examples/online_serving/elastic_ep/serve_deepseek_v2.sh:SC2086
-examples/online_serving/multi-node-serving.sh:SC2006
-examples/online_serving/multi-node-serving.sh:SC2086
-examples/online_serving/multi-node-serving.sh:SC2181
-examples/others/lmcache/disagg_prefill_lmcache_v1/disagg_example_nixl.sh:SC2046
-examples/others/lmcache/disagg_prefill_lmcache_v1/disagg_example_nixl.sh:SC2126
-examples/others/lmcache/disagg_prefill_lmcache_v1/disagg_example_nixl.sh:SC2181
-examples/others/lmcache/disagg_prefill_lmcache_v1/disagg_example_nixl.sh:SC2206
-examples/others/lmcache/disagg_prefill_lmcache_v1/disagg_vllm_launcher.sh:SC2086
-examples/pooling/embed/openai_embedding_long_text/service.sh:SC2086
-tests/standalone_tests/python_only_compile.sh:SC2086
-tests/v1/ec_connector/integration/run_epd_correctness_test.sh:SC2086
-tests/v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh:SC2086
-tests/v1/kv_connector/nixl_integration/run_accuracy_test.sh:SC2005
-tests/v1/kv_connector/nixl_integration/run_accuracy_test.sh:SC2086
-tests/v1/kv_connector/nixl_integration/run_accuracy_test.sh:SC2124
-tests/v1/kv_connector/nixl_integration/run_accuracy_test.sh:SC2126
-tests/v1/kv_connector/nixl_integration/run_accuracy_test.sh:SC2206
-tests/v1/kv_connector/nixl_integration/run_edge_case_test.sh:SC2086
-tests/v1/kv_connector/nixl_integration/run_edge_case_test.sh:SC2153
-tests/v1/kv_connector/nixl_integration/run_tpu_disagg_accuracy_test.sh:SC2086
-tests/v1/kv_connector/nixl_integration/run_tpu_disagg_accuracy_test.sh:SC2089
-tests/v1/kv_connector/nixl_integration/run_tpu_disagg_accuracy_test.sh:SC2090
-tests/v1/kv_connector/nixl_integration/run_tpu_edge_case_test.sh:SC2086
-tests/v1/kv_connector/nixl_integration/run_tpu_edge_case_test.sh:SC2089
-tests/v1/kv_connector/nixl_integration/run_tpu_edge_case_test.sh:SC2090
-tools/ep_kernels/elastic_ep/install_eep_libraries.sh:SC2086
-tools/ep_kernels/install_python_libraries.sh:SC2086
-tools/ep_kernels/install_python_libraries.sh:SC2196
-tools/flashinfer-build.sh:SC2086
-tools/flashinfer-build.sh:SC2269
-tools/install_deepgemm.sh:SC2035
-tools/install_deepgemm.sh:SC2295
-tools/pre_commit/shellcheck.sh:SC2016
-tools/vllm-rocm/generate-rocm-wheels-root-index.sh:SC2295
-tools/vllm-tpu/build.sh:SC2145

tools/pre_commit/shellcheck.sh

@@ -1,8 +1,7 @@
 #!/bin/bash
-set -euo pipefail
+set -e
 
 scversion="stable"
-baseline="tools/pre_commit/shellcheck.baseline"
 
 if [ -d "shellcheck-${scversion}" ]; then
     export PATH="$PATH:$(pwd)/shellcheck-${scversion}"
@@ -20,38 +19,4 @@ if ! [ -x "$(command -v shellcheck)" ]; then
 fi
 
 # TODO - fix warnings in .buildkite/scripts/hardware_ci/run-amd-test.sh
-# collects warnings as "file:SCcode" pairs for baseline comparison.
-collect() {
-  find . -path ./.git -prune -o -name "*.sh" \
-    -not -path "./.buildkite/scripts/hardware_ci/run-amd-test.sh" -print0 | \
-    xargs -0 sh -c 'for f in "$@"; do git check-ignore -q "$f" || shellcheck -s bash -f gcc "$f" || true; done' -- | \
-    sed -nE 's|^\./||; s|^([^:]+):[0-9]+:[0-9]+:.*\[(SC[0-9]+)\]$|\1:\2|p' | \
-    sort -u
-}
-
-if [[ "${1:-}" == "--generate-baseline" ]]; then
-  collect > "$baseline"
-  echo "Wrote baseline to $baseline"
-  exit 0
-fi
-
-if [[ ! -f "$baseline" ]]; then
-  echo "Baseline not found: $baseline (run: $0 --generate-baseline)"
-  exit 1
-fi
-
-current="$(mktemp)"
-trap 'rm -f "$current"' EXIT
-collect > "$current"
-
-# finds new warnings not in baseline
-new_errors="$(comm -23 "$current" <(sort -u "$baseline") || true)"
-if [ -n "$new_errors" ]; then
-  echo "$new_errors" | cut -d: -f1 | sort -u | while IFS= read -r file; do
-    if [[ -f "$file" ]]; then
-      codes=$(echo "$new_errors" | awk -F: -v f="$file" '$1==f {print $2}' | paste -sd ',' -)
-      shellcheck -s bash --include="$codes" "$file" 2>&1 || true
-    fi
-  done
-  exit 1
-fi
+find . -name "*.sh" ".git" -prune -not -path "./.buildkite/scripts/hardware_ci/run-amd-test.sh" -print0 | xargs -0 -I {} sh -c 'git check-ignore -q "{}" || shellcheck -s bash "{}"'

vllm/_custom_ops.py

@@ -3078,7 +3078,6 @@ def cpu_fused_moe(
     topk_ids: torch.Tensor,
     act: str,
     isa: str,
-    skip_weighted: bool = False,
 ) -> torch.Tensor:
     output = torch.empty_like(input)
     torch.ops._C.cpu_fused_moe(
@@ -3090,7 +3089,6 @@ def cpu_fused_moe(
         w2_bias,
         topk_weights,
         topk_ids,
-        skip_weighted,
         act,
         isa,
     )

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/lib/ready_checker.py

@@ -66,8 +66,7 @@ async def wait_for_endpoint(
                     pbar.close()
                     return output
                 else:
-                    err_last_line = str(output.error).rstrip().rsplit("\n", 1)[-1]
-                    logger.warning("Endpoint is not ready. Error='%s'", err_last_line)
+                    logger.warning("Endpoint is not ready. Error='%s'", output.error)
             except aiohttp.ClientConnectorError:
                 pass
 

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/model.py

@@ -1557,7 +1557,6 @@ class ModelConfig:
 
     @property
     def attn_type(self) -> AttnTypeStr:
-        """Determine the attention type based on model configuration."""
         if self.pooler_config is not None:
             seq_pooling_type = self._model_info.default_seq_pooling_type
             if seq_pooling_type == "CLS":

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,
@@ -278,7 +269,7 @@ class VllmConfig:
     optimization_level: OptimizationLevel = OptimizationLevel.O2
     """The optimization level. These levels trade startup time cost for
     performance, with -O0 having the best startup time and -O3 having the best
-    performance. -O2 is used by default. See OptimizationLevel for full
+    performance. -02 is used by defult. See  OptimizationLevel for full
     description."""
 
     weight_transfer_config: WeightTransferConfig | None = None

vllm/engine/async_llm_engine.py

@@ -4,4 +4,3 @@
 from vllm.v1.engine.async_llm import AsyncLLM
 
 AsyncLLMEngine = AsyncLLM  # type: ignore
-"""The `AsyncLLMEngine` class is an alias of [vllm.v1.engine.async_llm.AsyncLLM][]."""

vllm/engine/llm_engine.py

@@ -4,4 +4,3 @@
 from vllm.v1.engine.llm_engine import LLMEngine as V1LLMEngine
 
 LLMEngine = V1LLMEngine  # type: ignore
-"""The `LLMEngine` class is an alias of [vllm.v1.engine.llm_engine.LLMEngine][]."""

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/openai/realtime/connection.py

@@ -48,6 +48,7 @@ class RealtimeConnection:
         self.generation_task: asyncio.Task | None = None
 
         self._is_connected = False
+        self._is_input_finished = False
         self._is_model_validated = False
 
         self._max_audio_filesize_mb = envs.VLLM_MAX_AUDIO_CLIP_FILESIZE_MB
@@ -144,7 +145,7 @@ class RealtimeConnection:
             commit_event = InputAudioBufferCommit(**event)
             # final signals that the audio is finished
             if commit_event.final:
-                self.audio_queue.put_nowait(None)
+                self._is_input_finished = True
             else:
                 await self.start_generation()
         else:
@@ -238,6 +239,11 @@ class RealtimeConnection:
                     # finish because websocket connection was killed
                     break
 
+                if self.audio_queue.empty() and self._is_input_finished:
+                    # finish because client signals that audio input
+                    # is finished
+                    break
+
             usage = UsageInfo(
                 prompt_tokens=prompt_token_ids_len,
                 completion_tokens=completion_tokens_len,

vllm/entrypoints/openai/responses/serving.py

@@ -980,9 +980,7 @@ class OpenAIServingResponses(OpenAIServing):
             output_items.extend(last_items)
         return output_items
 
-    def _extract_system_message_from_request(
-        self, request: ResponsesRequest
-    ) -> str | None:
+    def _extract_system_message_from_request(self, request) -> str | None:
         system_msg = None
         if not isinstance(request.input, str):
             for response_msg in request.input:
@@ -990,17 +988,7 @@ class OpenAIServingResponses(OpenAIServing):
                     isinstance(response_msg, dict)
                     and response_msg.get("role") == "system"
                 ):
-                    content = response_msg.get("content")
-                    if isinstance(content, str):
-                        system_msg = content
-                    elif isinstance(content, list):
-                        for param in content:
-                            if (
-                                isinstance(param, dict)
-                                and param.get("type") == "input_text"
-                            ):
-                                system_msg = param.get("text")
-                                break
+                    system_msg = response_msg.get("content")
                     break
         return system_msg
 

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/forward_context.py

@@ -5,7 +5,7 @@ import time
 from collections import defaultdict
 from contextlib import contextmanager
 from dataclasses import dataclass, field
-from typing import Any
+from typing import Any, NamedTuple
 
 import torch
 
@@ -26,8 +26,7 @@ batchsize_logging_interval: float = envs.VLLM_LOG_BATCHSIZE_INTERVAL
 batchsize_forward_time: defaultdict = defaultdict(list)
 
 
-@dataclass(frozen=True)
-class BatchDescriptor:
+class BatchDescriptor(NamedTuple):
     """
     Batch descriptor for cudagraph dispatching. We should keep the num of
     items as minimal as possible to properly and uniquely describe the padded
@@ -57,6 +56,19 @@ class BatchDescriptor:
     to be properly captured.
     """
 
+    def relax_for_mixed_batch_cudagraphs(self) -> "BatchDescriptor":
+        """
+        Return a relaxed version of current batch descriptor that is still compatible
+        with PIECEWISE cudagraphs (or mixed prefill-decode FA cudagraphs).
+        """
+        return BatchDescriptor(
+            self.num_tokens,
+            num_reqs=None,
+            uniform=False,
+            has_lora=self.has_lora,
+            num_active_loras=self.num_active_loras,
+        )
+
 
 def _compute_sp_num_tokens(
     num_tokens_across_dp_cpu: torch.Tensor, sequence_parallel_size: int

vllm/inputs/data.py

@@ -298,7 +298,6 @@ which can be passed to
 
 
 SingletonInputs: TypeAlias = DecoderOnlyInputs | MultiModalEncDecInputs
-"""The inputs for a single encoder/decoder prompt."""
 
 
 @dataclass

vllm/model_executor/layers/fused_moe/__init__.py

@@ -4,11 +4,6 @@
 from contextlib import contextmanager
 from typing import Any
 
-from vllm.model_executor.layers.fused_moe.activation import (
-    MoEActivation,
-    activation_without_mul,
-    apply_moe_activation,
-)
 from vllm.model_executor.layers.fused_moe.config import (
     FusedMoEConfig,
     RoutingMethodType,
@@ -32,6 +27,7 @@ from vllm.model_executor.layers.fused_moe.shared_fused_moe import SharedFusedMoE
 from vllm.model_executor.layers.fused_moe.unquantized_fused_moe_method import (
     UnquantizedFusedMoEMethod,
 )
+from vllm.model_executor.layers.fused_moe.utils import activation_without_mul
 from vllm.model_executor.layers.fused_moe.zero_expert_fused_moe import (
     ZeroExpertFusedMoE,
 )
@@ -58,7 +54,6 @@ __all__ = [
     "FusedMoERouter",
     "FusedMoEConfig",
     "FusedMoEMethodBase",
-    "MoEActivation",
     "UnquantizedFusedMoEMethod",
     "FusedMoeWeightScaleSupported",
     "FusedMoEPermuteExpertsUnpermute",
@@ -68,7 +63,6 @@ __all__ = [
     "SharedFusedMoE",
     "ZeroExpertFusedMoE",
     "activation_without_mul",
-    "apply_moe_activation",
     "override_config",
     "get_config",
 ]
@@ -108,7 +102,6 @@ if HAS_TRITON:
     )
     from vllm.model_executor.layers.fused_moe.xpu_fused_moe import (
         XPUExperts,
-        XPUExpertsFp8,
     )
 
     __all__ += [
@@ -128,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/activation.py

@@ -1,136 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-"""MoE activation function enum and utilities."""
-
-from enum import Enum
-
-import torch
-import torch.nn.functional as F
-
-
-class MoEActivation(Enum):
-    """Activation functions for MoE layers."""
-
-    # Gated activations (gate * activation(up)) expect input of shape [..., 2*d]
-    # and produce output of shape [..., d]
-    SILU = "silu"
-    GELU = "gelu"
-    RELU2 = "relu2"
-    SWIGLUOAI = "swigluoai"
-    SWIGLUSTEP = "swiglustep"
-
-    # Non-gated activations (no mul with gate) expect input of shape [..., d]
-    # and produce output of shape [..., d].
-    # NOTE: Non-gated activations require the "_no_mul" suffix to be present.
-    SILU_NO_MUL = "silu_no_mul"
-    GELU_NO_MUL = "gelu_no_mul"
-    RELU2_NO_MUL = "relu2_no_mul"
-
-    @property
-    def is_gated(self) -> bool:
-        """Returns True if activation expects gate*activation(up) pattern.
-
-        Gated activations expect input tensor with 2x the output size,
-        where the first half is the gate and second half is the up projection.
-        """
-        return not self.value.endswith("_no_mul")
-
-    @property
-    def custom_op_name(self) -> str:
-        """Maps to the CustomOp name of activations
-        in vllm/model_executor/layers/activation.py."""
-        return _CUSTOM_OP_NAMES[self]
-
-    def without_mul(self) -> "MoEActivation":
-        """Get the non-gated variant of this activation.
-
-        For activations that have a _no_mul variant, returns that variant.
-        For activations without a _no_mul variant (or already _no_mul),
-        returns self.
-        """
-        return _WITHOUT_MUL.get(self, self)
-
-    @classmethod
-    def from_str(cls, s: str) -> "MoEActivation":
-        """Parse from string for backward compatibility."""
-        for member in cls:
-            if member.value == s:
-                return member
-        valid = [m.value for m in cls]
-        raise ValueError(f"Unknown MoE activation: {s!r}. Valid activations: {valid}")
-
-
-# Module-level lookup tables used by MoEActivation functions.
-_CUSTOM_OP_NAMES: dict[MoEActivation, str] = {
-    MoEActivation.SILU: "silu_and_mul",
-    MoEActivation.GELU: "gelu_and_mul",
-    MoEActivation.SWIGLUOAI: "swigluoai_and_mul",
-    MoEActivation.SWIGLUSTEP: "swiglustep_and_mul",
-    MoEActivation.RELU2: "relu2",
-    MoEActivation.SILU_NO_MUL: "silu_and_mul",
-    MoEActivation.GELU_NO_MUL: "gelu_and_mul",
-    MoEActivation.RELU2_NO_MUL: "relu2",
-}
-
-_WITHOUT_MUL: dict[MoEActivation, MoEActivation] = {
-    MoEActivation.SILU: MoEActivation.SILU_NO_MUL,
-    MoEActivation.GELU: MoEActivation.GELU_NO_MUL,
-    MoEActivation.RELU2: MoEActivation.RELU2_NO_MUL,
-}
-
-
-def activation_without_mul(activation: str) -> str:
-    """Get the non-gated variant of an activation function.
-
-    Args:
-        activation: The activation function name (e.g., "silu", "gelu")
-
-    Returns:
-        The non-gated activation name (e.g., "silu_no_mul", "gelu_no_mul")
-    """
-    return MoEActivation.from_str(activation).without_mul().value
-
-
-def apply_moe_activation(
-    activation: MoEActivation,
-    output: torch.Tensor,
-    input: torch.Tensor,
-) -> torch.Tensor:
-    """Apply MoE activation function."""
-    assert input.dim() == 2, "Input must be 2D"
-    assert output.dim() == 2, "Output must be 2D"
-    if activation.is_gated:
-        assert output.size(-1) * 2 == input.size(-1), (
-            f"{activation.value} expects 2x ratio: "
-            f"{output.size(-1) * 2} vs {input.size(-1)}"
-        )
-    else:
-        assert output.size(-1) == input.size(-1), (
-            f"{activation.value} expects equal sizes: "
-            f"{output.size(-1)} vs {input.size(-1)}"
-        )
-
-    # Activations with gated multiplication (gate × activation(up))
-    if activation == MoEActivation.SILU:
-        torch.ops._C.silu_and_mul(output, input)
-    elif activation == MoEActivation.GELU:
-        torch.ops._C.gelu_and_mul(output, input)
-    elif activation == MoEActivation.SWIGLUOAI:
-        torch.ops._C.swigluoai_and_mul(output, input)
-    elif activation == MoEActivation.SWIGLUSTEP:
-        from vllm.model_executor.layers.activation import swiglustep_and_mul_triton
-
-        swiglustep_and_mul_triton(output, input)
-
-    # Activations without gated multiplication
-    elif activation == MoEActivation.SILU_NO_MUL:
-        output.copy_(F.silu(input))
-    elif activation == MoEActivation.GELU_NO_MUL:
-        output.copy_(F.gelu(input))
-    elif activation == MoEActivation.RELU2_NO_MUL:
-        F.relu(input, inplace=True)
-        torch.square(input, out=output)
-    else:
-        raise ValueError(f"Unsupported FusedMoe activation: {activation}")
-
-    return output

vllm/model_executor/layers/fused_moe/batched_deep_gemm_moe.py

@@ -7,7 +7,6 @@ import torch
 import vllm.model_executor.layers.fused_moe.modular_kernel as mk
 from vllm.forward_context import get_forward_context, is_forward_context_available
 from vllm.logger import init_logger
-from vllm.model_executor.layers.fused_moe.activation import MoEActivation
 from vllm.model_executor.layers.fused_moe.config import (
     FusedMoEConfig,
     FusedMoEParallelConfig,
@@ -304,8 +303,8 @@ class BatchedDeepGemmExperts(mk.FusedMoEPermuteExpertsUnpermute):
         return (weight_key, activation_key) in SUPPORTED_W_A
 
     @staticmethod
-    def _supports_activation(activation: MoEActivation) -> bool:
-        return activation == MoEActivation.SILU
+    def _supports_activation(activation: str) -> bool:
+        return activation in ["silu"]
 
     @staticmethod
     def _supports_parallel_config(moe_parallel_config: FusedMoEParallelConfig) -> bool:
@@ -339,7 +338,7 @@ class BatchedDeepGemmExperts(mk.FusedMoEPermuteExpertsUnpermute):
         global_num_experts: int,
         local_num_experts: int,
         expert_tokens_meta: mk.ExpertTokensMetadata | None,
-        activation: MoEActivation,
+        activation: str,
     ) -> tuple[tuple[int, ...], tuple[int, ...], tuple[int, ...]]:
         # FIXME (varun): We should be able to dispatch only from the leader
         # DP ranks in the case of TP > 1. At the moment, all the Ranks
@@ -390,7 +389,7 @@ class BatchedDeepGemmExperts(mk.FusedMoEPermuteExpertsUnpermute):
         w2: torch.Tensor,
         topk_weights: torch.Tensor,
         topk_ids: torch.Tensor,
-        activation: MoEActivation,
+        activation: str,
         global_num_experts: int,
         expert_map: torch.Tensor | None,
         a1q_scale: torch.Tensor | None,

vllm/model_executor/layers/fused_moe/config.py

@@ -14,7 +14,6 @@ from vllm.distributed import (
     get_tensor_model_parallel_rank,
 )
 from vllm.logger import init_logger
-from vllm.model_executor.layers.fused_moe.activation import MoEActivation
 from vllm.model_executor.layers.quantization.utils.ocp_mx_utils import (
     OCP_MX_DTYPES,
     OCP_MX_Scheme,
@@ -1133,7 +1132,7 @@ class FusedMoEConfig:
     intermediate_size_per_partition: int
     num_local_experts: int
     num_logical_experts: int
-    activation: MoEActivation
+    activation: str
     device: torch.device | str
     routing_method: RoutingMethodType
     moe_parallel_config: FusedMoEParallelConfig

vllm/model_executor/layers/fused_moe/cpu_fused_moe.py

@@ -9,7 +9,6 @@ from torch.nn import functional as F
 from vllm import _custom_ops as ops
 from vllm._custom_ops import cpu_fused_moe, cpu_prepack_moe_weight
 from vllm.model_executor.layers.activation import SiluAndMul
-from vllm.model_executor.layers.fused_moe.activation import MoEActivation
 from vllm.model_executor.layers.quantization.utils.layer_utils import replace_parameter
 from vllm.utils.torch_utils import direct_register_custom_op
 
@@ -37,9 +36,9 @@ def _swigluoai_forward_native(
 # Map activation names to their native forward functions.
 # Uses static methods or standalone functions to avoid instantiating CustomOp
 # classes, which would call get_current_vllm_config() before config is set.
-_CPU_MOE_ACT_FN: dict[MoEActivation, Callable[[torch.Tensor], torch.Tensor]] = {
-    MoEActivation.SILU: SiluAndMul.forward_native,
-    MoEActivation.SWIGLUOAI: _swigluoai_forward_native,
+_CPU_MOE_ACT_FN: dict[str, Callable[[torch.Tensor], torch.Tensor]] = {
+    "silu": SiluAndMul.forward_native,
+    "swigluoai": _swigluoai_forward_native,
 }
 
 
@@ -169,9 +168,9 @@ class SGLFusedMOE:
         routed_scaling_factor: float = 1.0,
         e_score_correction_bias: torch.Tensor | None = None,
         apply_router_weight_on_input: bool = False,
-        activation: MoEActivation = MoEActivation.SILU,
+        activation: str = "silu",
     ) -> torch.Tensor:
-        assert activation == MoEActivation.SILU, f"{activation} is not supported."
+        assert activation == "silu", f"{activation} is not supported."
         assert not apply_router_weight_on_input
         topk_weights, topk_ids = select_experts(
             hidden_states=x,
@@ -207,8 +206,6 @@ class SGLFusedMOE:
 
 
 class CPUFusedMOE:
-    """CPU-based fused MoE implementation."""
-
     def __init__(self, layer: torch.nn.Module) -> None:
         use_grouped_gemm, isa = self.check_grouped_gemm(layer)
         self.isa = isa
@@ -236,9 +233,10 @@ class CPUFusedMOE:
         routed_scaling_factor: float = 1.0,
         e_score_correction_bias: torch.Tensor | None = None,
         apply_router_weight_on_input: bool = False,
-        activation: MoEActivation = MoEActivation.SILU,
+        activation: str = "silu",
     ) -> torch.Tensor:
         assert activation in _CPU_MOE_ACT_FN, f"{activation} is not supported."
+        assert not apply_router_weight_on_input
 
         topk_weights, topk_ids = select_experts(
             hidden_states=x,
@@ -261,7 +259,6 @@ class CPUFusedMOE:
             topk_ids,
             activation,
             global_num_experts,
-            apply_router_weight_on_input,
         )
 
     def check_grouped_gemm(
@@ -354,16 +351,9 @@ class CPUFusedMOE:
         input: torch.Tensor,
         topk_weights: torch.Tensor,
         topk_ids: torch.Tensor,
-        activation: MoEActivation,
+        activation: str,
         global_num_experts: int = -1,
-        skip_weighted: bool = False,
     ) -> torch.Tensor:
-        if skip_weighted:
-            assert topk_ids.size(1) == 1, (
-                "apply_router_weight_on_input is only implemented for topk=1"
-            )
-            input.mul_(topk_weights.to(input.dtype))
-
         output = cpu_fused_moe(
             input,
             layer.w13_weight,
@@ -372,9 +362,8 @@ class CPUFusedMOE:
             getattr(layer, "w2_bias", None),
             topk_weights,
             topk_ids,
-            activation.value,
+            activation,
             self.isa,
-            skip_weighted,
         )
         return output
 
@@ -384,16 +373,9 @@ class CPUFusedMOE:
         input: torch.Tensor,
         topk_weights: torch.Tensor,
         topk_ids: torch.Tensor,
-        activation: MoEActivation,
+        activation: str,
         global_num_experts: int = -1,
-        skip_weighted: bool = False,
     ) -> torch.Tensor:
-        if skip_weighted:
-            assert topk_ids.size(1) == 1, (
-                "apply_router_weight_on_input is only implemented for topk=1"
-            )
-            input.mul_(topk_weights.to(input.dtype))
-
         output = torch.empty_like(input)
         layer_id = id(layer)
         torch.ops.vllm.cpu_fused_moe_torch(
@@ -404,7 +386,6 @@ class CPUFusedMOE:
             topk_ids,
             activation,
             global_num_experts,
-            skip_weighted,
         )
 
         return output
@@ -418,9 +399,7 @@ def cpu_fused_moe_torch(
     topk_ids: torch.Tensor,
     activation: str,
     global_num_experts: int = -1,
-    skip_weighted: bool = False,
 ) -> None:
-    act = MoEActivation.from_str(activation)
     layer = _CPU_MOE_LAYER_CACHE[layer_id]()
 
     # Ref code from https://github.com/sgl-project/sglang/blob/716e682721397df103f347d22da8bd46c6016dab/python/sglang/srt/layers/moe/fused_moe_native.py#L53
@@ -444,7 +423,7 @@ def cpu_fused_moe_torch(
         tokens_for_this_expert = sorted_tokens[start_idx:end_idx]
 
         gate_up = layer.gate_up_linear[i](tokens_for_this_expert)  # type: ignore
-        gate_up = _CPU_MOE_ACT_FN[act](gate_up)
+        gate_up = _CPU_MOE_ACT_FN[activation](gate_up)
         expert_out = layer.down_linear[i](gate_up)  # type: ignore
         outputs.append(expert_out)
         start_idx = end_idx
@@ -453,16 +432,13 @@ def cpu_fused_moe_torch(
     new_x = torch.empty_like(outs)
 
     new_x[idxs] = outs
-    if skip_weighted:
-        final_out = new_x
-    else:
-        final_out = (
-            new_x.view(*topk_ids.shape, -1)
-            .type(topk_weights.dtype)
-            .mul_(topk_weights.unsqueeze(dim=-1))
-            .sum(dim=1)
-            .type(new_x.dtype)
-        )
+    final_out = (
+        new_x.view(*topk_ids.shape, -1)
+        .type(topk_weights.dtype)
+        .mul_(topk_weights.unsqueeze(dim=-1))
+        .sum(dim=1)
+        .type(new_x.dtype)
+    )
     output.copy_(final_out)
 
 

vllm/model_executor/layers/fused_moe/cutlass_moe.py

@@ -7,10 +7,6 @@ import torch
 import vllm.model_executor.layers.fused_moe.modular_kernel as mk
 from vllm import _custom_ops as ops
 from vllm.logger import init_logger
-from vllm.model_executor.layers.fused_moe.activation import (
-    MoEActivation,
-    apply_moe_activation,
-)
 from vllm.model_executor.layers.fused_moe.config import (
     FusedMoEConfig,
     FusedMoEParallelConfig,
@@ -29,6 +25,7 @@ from vllm.model_executor.layers.fused_moe.topk_weight_and_reduce import (
 )
 from vllm.model_executor.layers.fused_moe.utils import (
     _resize_cache,
+    apply_moe_activation,
 )
 from vllm.model_executor.layers.quantization.utils.quant_utils import (
     QuantKey,
@@ -54,7 +51,7 @@ def run_cutlass_moe_fp8(
     w1: torch.Tensor,
     w2: torch.Tensor,
     topk_ids: torch.Tensor,
-    activation: MoEActivation,
+    activation: str,
     global_num_experts: int,
     expert_map: torch.Tensor | None,
     w1_scale: torch.Tensor | None,
@@ -76,7 +73,7 @@ def run_cutlass_moe_fp8(
 ):
     a1q = hidden_states
 
-    assert activation.is_gated, "Only gated activation is supported"
+    assert not activation.endswith("_no_mul"), "Only gated activation is supported"
     assert w1_scale is not None
     assert w2_scale is not None
     assert w1.dtype == torch.float8_e4m3fn
@@ -313,12 +310,8 @@ class CutlassExpertsFp8Base(mk.FusedMoEPermuteExpertsUnpermute):
         return (weight_key, activation_key) in SUPPORTED_W_A
 
     @staticmethod
-    def _supports_activation(activation: MoEActivation) -> bool:
-        return activation in [
-            MoEActivation.SILU,
-            MoEActivation.GELU,
-            MoEActivation.SWIGLUOAI,
-        ]
+    def _supports_activation(activation: str) -> bool:
+        return activation in ["silu", "gelu", "swigluoai"]
 
     def finalize_weight_and_reduce_impl(self) -> mk.TopKWeightAndReduce:
         # Let PrepareAndFinalize::finalize() decide the impl.
@@ -332,7 +325,7 @@ class CutlassExpertsFp8Base(mk.FusedMoEPermuteExpertsUnpermute):
         w2: torch.Tensor,
         topk_weights: torch.Tensor,
         topk_ids: torch.Tensor,
-        activation: MoEActivation,
+        activation: str,
         global_num_experts: int,
         expert_map: torch.Tensor | None,
         a1q_scale: torch.Tensor | None,
@@ -383,8 +376,6 @@ class CutlassExpertsFp8Base(mk.FusedMoEPermuteExpertsUnpermute):
 
 
 class CutlassExpertsFp8(CutlassExpertsFp8Base):
-    """CUTLASS FP8 fused MoE expert implementation."""
-
     @staticmethod
     def activation_format() -> mk.FusedMoEActivationFormat:
         return mk.FusedMoEActivationFormat.Standard
@@ -422,7 +413,7 @@ class CutlassExpertsFp8(CutlassExpertsFp8Base):
         global_num_experts: int,
         local_num_experts: int,
         expert_tokens_meta: mk.ExpertTokensMetadata | None,
-        activation: MoEActivation,
+        activation: str,
     ) -> tuple[tuple[int, ...], tuple[int, ...], tuple[int, ...]]:
         activation_out_dim = self.adjust_N_for_activation(N, activation)
         workspace1 = (M * topk, max(N, K))
@@ -432,8 +423,6 @@ class CutlassExpertsFp8(CutlassExpertsFp8Base):
 
 
 class CutlassBatchedExpertsFp8(CutlassExpertsFp8Base):
-    """Batched CUTLASS FP8 fused MoE expert implementation."""
-
     @staticmethod
     def _supports_parallel_config(moe_parallel_config: FusedMoEParallelConfig) -> bool:
         # BATCHED activation format works with EP because
@@ -463,7 +452,7 @@ class CutlassBatchedExpertsFp8(CutlassExpertsFp8Base):
         global_num_experts: int,
         local_num_experts: int,
         expert_tokens_meta: mk.ExpertTokensMetadata | None,
-        activation: MoEActivation,
+        activation: str,
     ) -> tuple[tuple[int, ...], tuple[int, ...], tuple[int, ...]]:
         num_dp = self.num_dispatchers
         assert num_dp is not None
@@ -496,7 +485,7 @@ def run_cutlass_moe_fp4(
     w2_alphas: torch.Tensor,
     topk_weights: torch.Tensor,
     topk_ids: torch.Tensor,
-    activation: MoEActivation,
+    activation: str,
     workspace13: torch.Tensor,
     workspace2: torch.Tensor,
     m: int,
@@ -619,7 +608,7 @@ def run_cutlass_moe_fp4(
         blockscale_offsets[:-1],
     )
     del rep_a_fp4, rep_a_blockscale
-    if activation == MoEActivation.SILU:
+    if activation == "silu":
         # Fused SiLU+Mul+NVFP4 quantization
         # Note: c2 workspace is no longer needed since SiLU is fused with quantization.
         # c3 reuses workspace13 after c1 is consumed.
@@ -662,8 +651,6 @@ def run_cutlass_moe_fp4(
 
 
 class CutlassExpertsFp4(mk.FusedMoEPermuteExpertsUnpermute):
-    """CUTLASS FP4 fused MoE expert implementation."""
-
     @property
     def expects_unquantized_inputs(self) -> bool:
         return True
@@ -689,12 +676,8 @@ class CutlassExpertsFp4(mk.FusedMoEPermuteExpertsUnpermute):
         return (weight_key, activation_key) == (kNvfp4Static, kNvfp4Dynamic)
 
     @staticmethod
-    def _supports_activation(activation: MoEActivation) -> bool:
-        return activation in [
-            MoEActivation.SILU,
-            MoEActivation.GELU,
-            MoEActivation.SWIGLUOAI,
-        ]
+    def _supports_activation(activation: str) -> bool:
+        return activation in ["silu", "gelu", "swigluoai"]
 
     @staticmethod
     def _supports_parallel_config(moe_parallel_config: FusedMoEParallelConfig) -> bool:
@@ -727,7 +710,7 @@ class CutlassExpertsFp4(mk.FusedMoEPermuteExpertsUnpermute):
         global_num_experts: int,
         local_num_experts: int,
         expert_tokens_meta: mk.ExpertTokensMetadata | None,
-        activation: MoEActivation,
+        activation: str,
     ) -> tuple[tuple[int, ...], tuple[int, ...], tuple[int, ...]]:
         workspace1 = (M * topk, max(2 * N, K))
         workspace2 = (M * topk, N)
@@ -742,7 +725,7 @@ class CutlassExpertsFp4(mk.FusedMoEPermuteExpertsUnpermute):
         w2: torch.Tensor,
         topk_weights: torch.Tensor,
         topk_ids: torch.Tensor,
-        activation: MoEActivation,
+        activation: str,
         global_num_experts: int,
         expert_map: torch.Tensor | None,
         a1q_scale: torch.Tensor | None,  # unused
@@ -787,7 +770,7 @@ def run_cutlass_moe_w4a8_fp8(
     w1: torch.Tensor,
     w2: torch.Tensor,
     topk_ids: torch.Tensor,
-    activation: MoEActivation,
+    activation: str,
     global_num_experts: int,
     expert_map: torch.Tensor | None,
     w1_scale: torch.Tensor | None,
@@ -981,7 +964,7 @@ class CutlassExpertsW4A8Fp8(mk.FusedMoEPermuteExpertsUnpermute):
         )
 
     @staticmethod
-    def _supports_activation(activation: MoEActivation) -> bool:
+    def _supports_activation(activation: str) -> bool:
         raise NotImplementedError(
             "CutlassExpertsW4A8Fp8 is not yet used by an Oracle. "
             "This method should not be called."
@@ -1016,7 +999,7 @@ class CutlassExpertsW4A8Fp8(mk.FusedMoEPermuteExpertsUnpermute):
         global_num_experts: int,
         local_num_experts: int,
         expert_tokens_meta: mk.ExpertTokensMetadata | None,
-        activation: MoEActivation,
+        activation: str,
     ) -> tuple[tuple[int, ...], tuple[int, ...], tuple[int, ...]]:
         activation_out_dim = self.adjust_N_for_activation(N, activation)
         workspace1 = (M * topk, max(N, K))
@@ -1032,7 +1015,7 @@ class CutlassExpertsW4A8Fp8(mk.FusedMoEPermuteExpertsUnpermute):
         w2: torch.Tensor,
         topk_weights: torch.Tensor,
         topk_ids: torch.Tensor,
-        activation: MoEActivation,
+        activation: str,
         global_num_experts: int,
         expert_map: torch.Tensor | None,
         a1q_scale: torch.Tensor | None,
@@ -1105,7 +1088,7 @@ def cutlass_moe_w4a8_fp8(
     s_strides2: torch.Tensor,
     quant_config: FusedMoEQuantConfig,
     moe_config: FusedMoEConfig,
-    activation: MoEActivation = MoEActivation.SILU,
+    activation: str = "silu",
     expert_map: torch.Tensor | None = None,
     apply_router_weight_on_input: bool = False,
     global_num_experts: int = -1,
@@ -1148,7 +1131,7 @@ def cutlass_moe_w4a8_fp8(
         dtype: torch.int64
     - per_act_token (Optional[bool]): Whether the scale is per-token or
                                       per-tensor.
-    - activation (MoEActivation): The activation function to use.
+    - activation (str): The activation function to use.
     - expert_map (Optional[torch.Tensor]): In the case of Expert parallel,
         every Rank is responsible for a subset of experts. expert_map is a
         mapping from global expert-id to local expert-id. When expert_map[i]

vllm/model_executor/layers/fused_moe/deep_gemm_moe.py

@@ -5,7 +5,6 @@ import torch
 
 import vllm.model_executor.layers.fused_moe.modular_kernel as mk
 from vllm.logger import init_logger
-from vllm.model_executor.layers.fused_moe.activation import MoEActivation
 from vllm.model_executor.layers.fused_moe.config import (
     FusedMoEConfig,
     FusedMoEParallelConfig,
@@ -114,8 +113,6 @@ def _valid_deep_gemm(
 
 
 class DeepGemmExperts(mk.FusedMoEPermuteExpertsUnpermute):
-    """DeepGemm-based fused MoE expert implementation."""
-
     def __init__(self, moe_config: FusedMoEConfig, quant_config: FusedMoEQuantConfig):
         super().__init__(moe_config=moe_config, quant_config=quant_config)
         assert quant_config.block_shape == get_mk_alignment_for_contiguous_layout()
@@ -146,8 +143,8 @@ class DeepGemmExperts(mk.FusedMoEPermuteExpertsUnpermute):
         return (weight_key, activation_key) in SUPPORTED_W_A
 
     @staticmethod
-    def _supports_activation(activation: MoEActivation) -> bool:
-        return activation in [MoEActivation.SILU, MoEActivation.SWIGLUSTEP]
+    def _supports_activation(activation: str) -> bool:
+        return activation in ["silu", "swiglustep"]
 
     @staticmethod
     def _supports_parallel_config(moe_parallel_config: FusedMoEParallelConfig) -> bool:
@@ -172,7 +169,7 @@ class DeepGemmExperts(mk.FusedMoEPermuteExpertsUnpermute):
         global_num_experts: int,
         local_num_experts: int,
         expert_tokens_meta: mk.ExpertTokensMetadata | None,
-        activation: MoEActivation,
+        activation: str,
     ) -> tuple[tuple[int, ...], tuple[int, ...], tuple[int, ...]]:
         assert self.block_shape is not None
         block_m = self.block_shape[0]
@@ -188,7 +185,7 @@ class DeepGemmExperts(mk.FusedMoEPermuteExpertsUnpermute):
         return (workspace1, workspace2, output)
 
     def _act_mul_quant(
-        self, input: torch.Tensor, output: torch.Tensor, activation: MoEActivation
+        self, input: torch.Tensor, output: torch.Tensor, activation: str
     ) -> tuple[torch.Tensor, torch.Tensor]:
         assert self.block_shape is not None
         block_k = self.block_shape[1]
@@ -211,7 +208,7 @@ class DeepGemmExperts(mk.FusedMoEPermuteExpertsUnpermute):
             return a2q, a2q_scale
 
         # 2. Hopper / non‑E8M0: prefer the fused SiLU+mul+quant kernel
-        if activation == MoEActivation.SILU:
+        if activation == "silu":
             use_ue8m0 = scale_fmt == DeepGemmQuantScaleFMT.FLOAT32_CEIL_UE8M0
             return silu_mul_per_token_group_quant_fp8_colmajor(
                 input=input,
@@ -236,7 +233,7 @@ class DeepGemmExperts(mk.FusedMoEPermuteExpertsUnpermute):
         w2: torch.Tensor,
         topk_weights: torch.Tensor,
         topk_ids: torch.Tensor,
-        activation: MoEActivation,
+        activation: str,
         global_num_experts: int,
         expert_map: torch.Tensor | None,
         a1q_scale: torch.Tensor | None,

vllm/model_executor/layers/fused_moe/fallback.py

@@ -6,7 +6,6 @@ from abc import ABC, abstractmethod
 import torch
 
 import vllm.model_executor.layers.fused_moe.modular_kernel as mk
-from vllm.model_executor.layers.fused_moe.activation import MoEActivation
 from vllm.model_executor.layers.fused_moe.config import FusedMoEParallelConfig
 from vllm.model_executor.layers.quantization.utils.quant_utils import QuantKey
 
@@ -77,7 +76,7 @@ class FallbackExperts(mk.FusedMoEPermuteExpertsUnpermute, ABC):
         ) and fallback_cls._supports_quant_scheme(weight_key, activation_key)
 
     @classmethod
-    def _supports_activation(cls, activation: MoEActivation) -> bool:
+    def _supports_activation(cls, activation: str) -> bool:
         experts_cls, fallback_cls = cls.get_clses()
         return experts_cls._supports_activation(
             activation
@@ -139,7 +138,7 @@ class FallbackExperts(mk.FusedMoEPermuteExpertsUnpermute, ABC):
         global_num_experts: int,
         local_num_experts: int,
         expert_tokens_meta: mk.ExpertTokensMetadata | None,
-        activation: MoEActivation,
+        activation: str,
     ) -> tuple[tuple[int, ...], tuple[int, ...], tuple[int, ...]]:
         raise NotImplementedError
 
@@ -160,7 +159,7 @@ class FallbackExperts(mk.FusedMoEPermuteExpertsUnpermute, ABC):
         w2: torch.Tensor,
         topk_weights: torch.Tensor,
         topk_ids: torch.Tensor,
-        activation: MoEActivation,
+        activation: str,
         global_num_experts: int,
         expert_map: torch.Tensor | None,
         a1q_scale: torch.Tensor | None,

vllm/model_executor/layers/fused_moe/flashinfer_cutedsl_moe.py

@@ -6,7 +6,6 @@ import torch
 import vllm.model_executor.layers.fused_moe.modular_kernel as mk
 from vllm import envs
 from vllm.logger import init_logger
-from vllm.model_executor.layers.fused_moe.activation import MoEActivation
 from vllm.model_executor.layers.fused_moe.config import (
     FusedMoEConfig,
     FusedMoEParallelConfig,
@@ -73,8 +72,8 @@ class FlashInferCuteDSLExperts(mk.FusedMoEPermuteExpertsUnpermute):
         return (weight_key, activation_key) in SUPPORTED_W_A
 
     @staticmethod
-    def _supports_activation(activation: MoEActivation) -> bool:
-        return activation == MoEActivation.SILU
+    def _supports_activation(activation: str) -> bool:
+        return activation in ["silu"]
 
     @staticmethod
     def _supports_parallel_config(moe_parallel_config: FusedMoEParallelConfig) -> bool:
@@ -102,7 +101,7 @@ class FlashInferCuteDSLExperts(mk.FusedMoEPermuteExpertsUnpermute):
         global_num_experts: int,
         local_num_experts: int,
         expert_tokens_meta: mk.ExpertTokensMetadata | None,
-        activation: MoEActivation,
+        activation: str,
     ) -> tuple[tuple[int, ...], tuple[int, ...], tuple[int, ...]]:
         # We use global_num_experts due to how moe_align_block_size handles
         # expert_maps.
@@ -136,7 +135,7 @@ class FlashInferCuteDSLExperts(mk.FusedMoEPermuteExpertsUnpermute):
         w2: torch.Tensor,
         topk_weights: torch.Tensor,
         topk_ids: torch.Tensor,
-        activation: MoEActivation,
+        activation: str,
         global_num_experts: int,
         expert_map: torch.Tensor | None,
         a1q_scale: torch.Tensor | None,

vllm/model_executor/layers/fused_moe/flashinfer_cutlass_moe.py

@@ -5,7 +5,6 @@ import torch
 
 import vllm.model_executor.layers.fused_moe.modular_kernel as mk
 from vllm.logger import init_logger
-from vllm.model_executor.layers.fused_moe.activation import MoEActivation
 from vllm.model_executor.layers.fused_moe.config import (
     FusedMoEParallelConfig,
     FusedMoEQuantConfig,
@@ -131,8 +130,8 @@ class FlashInferExperts(mk.FusedMoEPermuteExpertsUnpermute):
         )
 
     @staticmethod
-    def _supports_activation(activation: MoEActivation) -> bool:
-        return activation in [MoEActivation.SILU, MoEActivation.RELU2_NO_MUL]
+    def _supports_activation(activation: str) -> bool:
+        return activation in ["silu", "relu2_no_mul"]
 
     @staticmethod
     def _supports_parallel_config(moe_parallel_config: FusedMoEParallelConfig) -> bool:
@@ -165,7 +164,7 @@ class FlashInferExperts(mk.FusedMoEPermuteExpertsUnpermute):
         global_num_experts: int,
         local_num_experts: int,
         expert_tokens_meta: mk.ExpertTokensMetadata | None,
-        activation: MoEActivation,
+        activation: str,
     ) -> tuple[tuple[int, ...], tuple[int, ...], tuple[int, ...]]:
         # We use global_num_experts due to how moe_align_block_size handles
         # expert_maps.
@@ -202,7 +201,7 @@ class FlashInferExperts(mk.FusedMoEPermuteExpertsUnpermute):
         w2: torch.Tensor,
         topk_weights: torch.Tensor,
         topk_ids: torch.Tensor,
-        activation: MoEActivation,
+        activation: str,
         global_num_experts: int,
         expert_map: torch.Tensor | None,
         a1q_scale: torch.Tensor | None,
@@ -215,8 +214,8 @@ class FlashInferExperts(mk.FusedMoEPermuteExpertsUnpermute):
         from flashinfer.fused_moe.core import ActivationType
 
         activation_str_to_value_map = {
-            MoEActivation.SILU: ActivationType.Swiglu,  # This is the default
-            MoEActivation.RELU2_NO_MUL: ActivationType.Relu2,
+            "silu": ActivationType.Swiglu,  # This is the default
+            "relu2_no_mul": ActivationType.Relu2,
         }
         assert activation in activation_str_to_value_map, (
             f"{activation=} missing from {activation_str_to_value_map.keys()=}"

vllm/model_executor/layers/fused_moe/flashinfer_trtllm_moe.py

@@ -4,7 +4,6 @@
 import torch
 
 import vllm.model_executor.layers.fused_moe.modular_kernel as mk
-from vllm.model_executor.layers.fused_moe.activation import MoEActivation
 from vllm.model_executor.layers.fused_moe.config import (
     FusedMoEConfig,
     FusedMoEParallelConfig,
@@ -51,9 +50,9 @@ def _supports_quant_scheme(
     return (weight_key, activation_key) in SUPPORTED_W_A
 
 
-def _supports_activation(activation: MoEActivation) -> bool:
+def _supports_activation(activation: str) -> bool:
     """Supports silu activation only."""
-    return activation == MoEActivation.SILU
+    return activation in ["silu"]
 
 
 def _supports_routing_method(

vllm/model_executor/layers/fused_moe/fused_batched_moe.py

@@ -5,7 +5,6 @@
 import torch
 
 import vllm.model_executor.layers.fused_moe.modular_kernel as mk
-from vllm.model_executor.layers.fused_moe.activation import MoEActivation
 from vllm.model_executor.layers.fused_moe.config import (
     FusedMoEConfig,
     FusedMoEParallelConfig,
@@ -699,7 +698,7 @@ class NaiveBatchedExperts(mk.FusedMoEPermuteExpertsUnpermute):
         )
 
     @staticmethod
-    def _supports_activation(activation: MoEActivation) -> bool:
+    def _supports_activation(activation: str) -> bool:
         raise NotImplementedError(
             "NaiveBatchedExperts is not yet used by an Oracle. "
             "This method should not be called."
@@ -731,7 +730,7 @@ class NaiveBatchedExperts(mk.FusedMoEPermuteExpertsUnpermute):
         global_num_experts: int,
         local_num_experts: int,
         expert_tokens_meta: mk.ExpertTokensMetadata | None,
-        activation: MoEActivation,
+        activation: str,
     ) -> tuple[tuple[int, ...], tuple[int, ...], tuple[int, ...]]:
         assert self.num_dispatchers is not None
         assert self.max_num_tokens is not None
@@ -758,7 +757,7 @@ class NaiveBatchedExperts(mk.FusedMoEPermuteExpertsUnpermute):
         w2: torch.Tensor,
         topk_weights: torch.Tensor,
         topk_ids: torch.Tensor,
-        activation: MoEActivation,
+        activation: str,
         global_num_experts: int,
         expert_map: torch.Tensor | None,
         a1q_scale: torch.Tensor | None,
@@ -943,14 +942,14 @@ class BatchedTritonExperts(mk.FusedMoEPermuteExpertsUnpermute):
         )
 
     @staticmethod
-    def _supports_activation(activation: MoEActivation) -> bool:
+    def _supports_activation(activation: str) -> bool:
         return activation in [
-            MoEActivation.SILU,
-            MoEActivation.GELU,
-            MoEActivation.SWIGLUOAI,
-            MoEActivation.SILU_NO_MUL,
-            MoEActivation.GELU_NO_MUL,
-            MoEActivation.RELU2_NO_MUL,
+            "silu",
+            "gelu",
+            "swigluoai",
+            "silu_no_mul",
+            "gelu_no_mul",
+            "relu2_no_mul",
         ]
 
     @staticmethod
@@ -976,7 +975,7 @@ class BatchedTritonExperts(mk.FusedMoEPermuteExpertsUnpermute):
         global_num_experts: int,
         local_num_experts: int,
         expert_tokens_meta: mk.ExpertTokensMetadata | None,
-        activation: MoEActivation,
+        activation: str,
     ) -> tuple[tuple[int, ...], tuple[int, ...], tuple[int, ...]]:
         assert self.num_dispatchers is not None
         assert self.max_num_tokens is not None
@@ -997,7 +996,7 @@ class BatchedTritonExperts(mk.FusedMoEPermuteExpertsUnpermute):
         w2: torch.Tensor,
         topk_weights: torch.Tensor,
         topk_ids: torch.Tensor,
-        activation: MoEActivation,
+        activation: str,
         global_num_experts: int,
         expert_map: torch.Tensor | None,
         a1q_scale: torch.Tensor | None,

vllm/model_executor/layers/fused_moe/fused_marlin_moe.py

@@ -8,10 +8,6 @@ import torch
 
 import vllm._custom_ops as ops
 import vllm.model_executor.layers.fused_moe.modular_kernel as mk
-from vllm.model_executor.layers.fused_moe.activation import (
-    MoEActivation,
-    apply_moe_activation,
-)
 from vllm.model_executor.layers.fused_moe.config import (
     FusedMoEConfig,
     FusedMoEParallelConfig,
@@ -27,6 +23,7 @@ from vllm.model_executor.layers.fused_moe.topk_weight_and_reduce import (
 )
 from vllm.model_executor.layers.fused_moe.utils import (
     _resize_cache,
+    apply_moe_activation,
     disable_inplace,
 )
 from vllm.model_executor.layers.quantization.utils.marlin_utils import (
@@ -62,9 +59,9 @@ def _fused_marlin_moe(
     sorted_token_ids: torch.Tensor,
     expert_ids: torch.Tensor,
     num_tokens_post_padded: torch.Tensor,
-    activation: MoEActivation = MoEActivation.SILU,
+    activation: str = "silu",
     activation_func: Callable[
-        [MoEActivation, torch.Tensor, torch.Tensor], None
+        [str, torch.Tensor, torch.Tensor], None
     ] = apply_moe_activation,
     input_global_scale1: torch.Tensor | None = None,
     input_global_scale2: torch.Tensor | None = None,
@@ -86,7 +83,7 @@ def _fused_marlin_moe(
     assert hidden_states.ndim == 2
     M, K = hidden_states.size()
     N = marlin_moe_intermediate_size(w1, w2)
-    w13_num_shards = 2 if activation.is_gated else 1
+    w13_num_shards = 1 if "no_mul" in activation else 2
     if workspace is None:
         workspace = marlin_make_workspace_new(hidden_states.device, 4)
 
@@ -218,9 +215,9 @@ def fused_marlin_moe(
     quant_type_id: int,
     apply_router_weight_on_input: bool = False,
     global_num_experts: int = -1,
-    activation: MoEActivation = MoEActivation.SILU,
+    activation: str = "silu",
     activation_func: Callable[
-        [MoEActivation, torch.Tensor, torch.Tensor], None
+        [str, torch.Tensor, torch.Tensor], None
     ] = apply_moe_activation,
     moe_sum: Callable[[torch.Tensor, torch.Tensor], None] | None = None,
     expert_map: torch.Tensor | None = None,
@@ -380,7 +377,7 @@ def batched_fused_marlin_moe(
     quant_type_id: int,
     apply_router_weight_on_input: bool = False,
     global_num_experts: int = -1,
-    activation: MoEActivation = MoEActivation.SILU,
+    activation: str | None = "silu",
     expert_map: torch.Tensor | None = None,
     global_scale1: torch.Tensor | None = None,
     global_scale2: torch.Tensor | None = None,
@@ -582,14 +579,14 @@ class MarlinExpertsBase(mk.FusedMoEPermuteExpertsUnpermute):
         return weight_key in SUPPORTED_W
 
     @staticmethod
-    def _supports_activation(activation: MoEActivation) -> bool:
+    def _supports_activation(activation: str) -> bool:
         return activation in [
-            MoEActivation.SILU,
-            MoEActivation.GELU,
-            MoEActivation.SWIGLUOAI,
-            MoEActivation.SILU_NO_MUL,
-            MoEActivation.GELU_NO_MUL,
-            MoEActivation.RELU2_NO_MUL,
+            "silu",
+            "gelu",
+            "swigluoai",
+            "silu_no_mul",
+            "gelu_no_mul",
+            "relu2_no_mul",
         ]
 
     @staticmethod
@@ -640,8 +637,6 @@ class MarlinExpertsBase(mk.FusedMoEPermuteExpertsUnpermute):
 
 
 class MarlinExperts(MarlinExpertsBase):
-    """Marlin-based fused MoE expert implementation."""
-
     def supports_expert_map(self) -> bool:
         return True
 
@@ -664,7 +659,7 @@ class MarlinExperts(MarlinExpertsBase):
         global_num_experts: int,
         local_num_experts: int,
         expert_tokens_meta: mk.ExpertTokensMetadata | None,
-        activation: MoEActivation,
+        activation: str,
     ) -> tuple[tuple[int, ...], tuple[int, ...], tuple[int, ...]]:
         # Modular Kernel provisions output buffer from workspace1. However in
         # the fused_marlin_moe() function, the final torch.sum(), is defined
@@ -695,7 +690,7 @@ class MarlinExperts(MarlinExpertsBase):
         w2: torch.Tensor,
         topk_weights: torch.Tensor,
         topk_ids: torch.Tensor,
-        activation: MoEActivation,
+        activation: str,
         global_num_experts: int,
         expert_map: torch.Tensor | None,
         a1q_scale: torch.Tensor | None,
@@ -743,8 +738,6 @@ class MarlinExperts(MarlinExpertsBase):
 
 
 class BatchedMarlinExperts(MarlinExpertsBase):
-    """Batched Marlin-based fused MoE expert implementation."""
-
     def __init__(
         self,
         moe_config: FusedMoEConfig,
@@ -791,7 +784,7 @@ class BatchedMarlinExperts(MarlinExpertsBase):
         global_num_experts: int,
         local_num_experts: int,
         expert_tokens_meta: mk.ExpertTokensMetadata | None,
-        activation: MoEActivation,
+        activation: str,
     ) -> tuple[tuple[int, ...], tuple[int, ...], tuple[int, ...]]:
         assert self.num_dispatchers is not None
         assert self.max_num_tokens is not None
@@ -811,7 +804,7 @@ class BatchedMarlinExperts(MarlinExpertsBase):
         w2: torch.Tensor,
         topk_weights: torch.Tensor,
         topk_ids: torch.Tensor,
-        activation: MoEActivation,
+        activation: str,
         global_num_experts: int,
         expert_map: torch.Tensor | None,
         a1q_scale: torch.Tensor | None,

vllm/model_executor/layers/fused_moe/fused_moe.py

@@ -17,10 +17,6 @@ from vllm.logger import init_logger
 from vllm.model_executor.layers.batch_invariant import (
     vllm_is_batch_invariant,
 )
-from vllm.model_executor.layers.fused_moe.activation import (
-    MoEActivation,
-    apply_moe_activation,
-)
 from vllm.model_executor.layers.fused_moe.config import (
     FUSED_MOE_UNQUANTIZED_CONFIG,
     FusedMoEConfig,
@@ -36,6 +32,7 @@ from vllm.model_executor.layers.fused_moe.topk_weight_and_reduce import (
 )
 from vllm.model_executor.layers.fused_moe.utils import (
     _resize_cache,
+    apply_moe_activation,
     disable_inplace,
     moe_kernel_quantize_input,
 )
@@ -98,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
@@ -332,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,
@@ -1471,7 +1468,6 @@ def outplace_fused_experts_fake(
     topk_weights: torch.Tensor,
     topk_ids: torch.Tensor,
     activation: str = "silu",
-    apply_router_weight_on_input: bool = False,
     use_fp8_w8a8: bool = False,
     use_int8_w8a8: bool = False,
     use_int8_w8a16: bool = False,
@@ -1525,13 +1521,12 @@ def fused_experts(
     topk_weights: torch.Tensor,
     topk_ids: torch.Tensor,
     inplace: bool = False,
-    activation: MoEActivation = MoEActivation.SILU,
+    activation: str = "silu",
     apply_router_weight_on_input: bool = False,
     global_num_experts: int = -1,
     expert_map: torch.Tensor | None = None,
     quant_config: FusedMoEQuantConfig | None = None,
 ) -> torch.Tensor:
-    """Run fused MoE expert computation using Triton kernels."""
     if quant_config is None:
         quant_config = FUSED_MOE_UNQUANTIZED_CONFIG
 
@@ -1543,7 +1538,7 @@ def fused_experts(
         w2=w2,
         topk_weights=topk_weights,
         topk_ids=topk_ids,
-        activation=activation.value,
+        activation=activation,
         apply_router_weight_on_input=apply_router_weight_on_input,
         use_fp8_w8a8=quant_config.use_fp8_w8a8,
         use_int8_w8a8=quant_config.use_int8_w8a8,
@@ -1622,9 +1617,6 @@ def fused_experts_impl(
     w1_bias: torch.Tensor | None = None,
     w2_bias: torch.Tensor | None = None,
 ) -> torch.Tensor:
-    # Convert string activation to enum for internal use
-    activation_enum = MoEActivation.from_str(activation)
-
     # Check constraints.
     if use_int4_w4a16:
         assert hidden_states.size(1) // 2 == w1.size(2), "Hidden size mismatch"
@@ -1699,7 +1691,7 @@ def fused_experts_impl(
 
     # This needs separate memory since it's used concurrently with cache1
     activation_out_dim = mk.FusedMoEPermuteExpertsUnpermute.adjust_N_for_activation(
-        N, activation_enum
+        N, activation
     )
     intermediate_cache2 = torch.empty(
         (M * top_k_num, activation_out_dim),
@@ -1839,7 +1831,7 @@ def fused_experts_impl(
         )
 
         apply_moe_activation(
-            activation_enum, intermediate_cache2, intermediate_cache1.view(-1, N)
+            activation, intermediate_cache2, intermediate_cache1.view(-1, N)
         )
 
         qintermediate_cache2, a2q_scale = moe_kernel_quantize_input(
@@ -1887,8 +1879,6 @@ def fused_experts_impl(
 
 
 class TritonExperts(mk.FusedMoEPermuteExpertsUnpermute):
-    """Triton-based fused MoE expert implementation."""
-
     def __init__(
         self,
         moe_config: FusedMoEConfig,
@@ -1939,13 +1929,8 @@ class TritonExperts(mk.FusedMoEPermuteExpertsUnpermute):
         return (weight_key, activation_key) in SUPPORTED_W_A
 
     @staticmethod
-    def _supports_activation(activation: MoEActivation) -> bool:
-        return activation in [
-            MoEActivation.SILU,
-            MoEActivation.GELU,
-            MoEActivation.SWIGLUOAI,
-            MoEActivation.SWIGLUSTEP,
-        ]
+    def _supports_activation(activation: str) -> bool:
+        return activation in ["silu", "gelu", "swigluoai", "swiglustep"]
 
     @staticmethod
     def _supports_parallel_config(moe_parallel_config: FusedMoEParallelConfig) -> bool:
@@ -1969,7 +1954,7 @@ class TritonExperts(mk.FusedMoEPermuteExpertsUnpermute):
         global_num_experts: int,
         local_num_experts: int,
         expert_tokens_meta: mk.ExpertTokensMetadata | None,
-        activation: MoEActivation,
+        activation: str,
     ) -> tuple[tuple[int, ...], tuple[int, ...], tuple[int, ...]]:
         activation_out_dim = self.adjust_N_for_activation(N, activation)
         workspace1 = (M, topk, max(activation_out_dim, K))
@@ -1985,7 +1970,7 @@ class TritonExperts(mk.FusedMoEPermuteExpertsUnpermute):
         w2: torch.Tensor,
         topk_weights: torch.Tensor,
         topk_ids: torch.Tensor,
-        activation: MoEActivation,
+        activation: str,
         global_num_experts: int,
         expert_map: torch.Tensor | None,
         a1q_scale: torch.Tensor | None,
@@ -2150,7 +2135,7 @@ class TritonWNA16Experts(TritonExperts):
         )
 
     @staticmethod
-    def _supports_activation(activation: MoEActivation) -> bool:
+    def _supports_activation(activation: str) -> bool:
         raise NotImplementedError(
             "TritonWNA16Experts is not yet used by an Oracle. "
             "This method should not be called."
@@ -2171,7 +2156,7 @@ class TritonWNA16Experts(TritonExperts):
         w2: torch.Tensor,
         topk_weights: torch.Tensor,
         topk_ids: torch.Tensor,
-        activation: MoEActivation,
+        activation: str,
         global_num_experts: int,
         expert_map: torch.Tensor | None,
         a1q_scale: torch.Tensor | None,

vllm/model_executor/layers/fused_moe/gpt_oss_triton_kernels_moe.py

@@ -7,7 +7,6 @@ import torch
 import vllm.model_executor.layers.fused_moe.modular_kernel as mk
 from vllm import _custom_ops as ops
 from vllm.logger import init_logger
-from vllm.model_executor.layers.fused_moe.activation import MoEActivation
 from vllm.model_executor.layers.fused_moe.config import (
     FUSED_MOE_UNQUANTIZED_CONFIG,
     FusedMoEParallelConfig,
@@ -173,7 +172,7 @@ def triton_kernel_moe_forward(
     gating_output: torch.Tensor,
     topk: int,
     renormalize: bool,
-    activation: MoEActivation = MoEActivation.SWIGLUOAI,
+    activation: str = "silu",
     quant_config: FusedMoEQuantConfig | None = None,
     apply_router_weight_on_input: bool = False,
     global_num_experts: int = -1,
@@ -212,7 +211,7 @@ def triton_kernel_fused_experts(
     gather_indx,  # GatherIndx
     scatter_indx,  # ScatterIndx
     topk: int,
-    activation: MoEActivation = MoEActivation.SWIGLUOAI,
+    activation: str = "silu",
     quant_config: FusedMoEQuantConfig | None = None,
     swiglu_alpha: float = 1.702,
     swiglu_limit: float = 7.0,
@@ -222,10 +221,6 @@ def triton_kernel_fused_experts(
     intermediate_cache: torch.Tensor | None = None,
     a1q_scale: torch.Tensor | None = None,
 ) -> torch.Tensor:
-    """Triton implementation of fused expert computation using OAI kernels."""
-    assert activation == MoEActivation.SWIGLUOAI, (
-        "Only SWIGLUOAI activation is supported"
-    )
     if quant_config is None:
         quant_config = FUSED_MOE_UNQUANTIZED_CONFIG
 
@@ -383,7 +378,7 @@ class BaseOAITritonExperts(mk.FusedMoEPermuteExpertsUnpermute):
         )
 
     @staticmethod
-    def _supports_activation(activation: MoEActivation) -> bool:
+    def _supports_activation(activation: str) -> bool:
         raise NotImplementedError(
             "OAITritonExperts is not yet used by an Oracle. "
             "This method should not be called."
@@ -449,8 +444,6 @@ class BaseOAITritonExperts(mk.FusedMoEPermuteExpertsUnpermute):
 
 
 class OAITritonExperts(BaseOAITritonExperts):
-    """OAI Triton-based fused MoE expert implementation."""
-
     @staticmethod
     def activation_format() -> mk.FusedMoEActivationFormat:
         return mk.FusedMoEActivationFormat.Standard
@@ -467,7 +460,7 @@ class OAITritonExperts(BaseOAITritonExperts):
         global_num_experts: int,
         local_num_experts: int,
         expert_tokens_meta: mk.ExpertTokensMetadata | None,
-        activation: MoEActivation,
+        activation: str,
     ) -> tuple[tuple[int, ...], tuple[int, ...], tuple[int, ...]]:
         # workspace are allocated inside the kernel
         activation_out_dim = self.adjust_N_for_activation(N, activation)
@@ -484,7 +477,7 @@ class OAITritonExperts(BaseOAITritonExperts):
         w2: torch.Tensor,
         topk_weights: torch.Tensor,
         topk_ids: torch.Tensor,
-        activation: MoEActivation,
+        activation: str,
         global_num_experts: int,
         expert_map: torch.Tensor | None,
         a1q_scale: torch.Tensor | None,
@@ -551,7 +544,7 @@ class UnfusedOAITritonExperts(BaseOAITritonExperts):
         global_num_experts: int,
         local_num_experts: int,
         expert_tokens_meta: mk.ExpertTokensMetadata | None,
-        activation: MoEActivation,
+        activation: str,
     ) -> tuple[tuple[int, ...], tuple[int, ...], tuple[int, ...]]:
         # workspace are allocated inside the kernel
         activation_out_dim = self.adjust_N_for_activation(N, activation)
@@ -571,7 +564,7 @@ class UnfusedOAITritonExperts(BaseOAITritonExperts):
         w2: torch.Tensor,
         topk_weights: torch.Tensor,
         topk_ids: torch.Tensor,
-        activation: MoEActivation,
+        activation: str,
         global_num_experts: int,
         expert_map: torch.Tensor | None,
         a1q_scale: torch.Tensor | None,