A high-throughput and memory-efficient inference and serving engine for LLMs
Mirrored from https://github.com/tduncan-lum/Nvidia-Optimization-Stack---VLLM.git
Synced Nov 3, 20:15 UTC
[MoE Refactor] Move remaining experts classes to experts directory (#42334)
Signed-off-by: Bill Nell <bnell@redhat.com>
bnellnm committed
May 12, 2026 at 09:19 UTC
6427603ae81353e34d0a9e066674e510c9d840ef
26 files changed
+48
-42
benchmarks/kernels/benchmark_moe.py
+2
-2
@@ -27,10 +27,10 @@ from vllm.model_executor.layers.fused_moe.config import (
27
RoutingMethodType,
28
_get_config_dtype_str,
29
)
30
-from vllm.model_executor.layers.fused_moe.fused_moe import *
31
-from vllm.model_executor.layers.fused_moe.triton_deep_gemm_moe import (
30
+from vllm.model_executor.layers.fused_moe.experts.triton_deep_gemm_moe import (
31
TritonOrDeepGemmExperts,
32
)
33
+from vllm.model_executor.layers.fused_moe.fused_moe import *
34
from vllm.transformers_utils.config import get_config
35
from vllm.triton_utils import triton
36
from vllm.utils.argparse_utils import FlexibleArgumentParser
docs/design/moe_kernel_features.md
+2
-2
@@ -81,7 +81,7 @@ To be used with a particular `FusedMoEPrepareAndFinalizeModular` subclass, MoE k
81
| Kernel | Input act. format | Quant. types | Quant. format | Activation function | Apply Weight On Input | Modular | Source |
82
| ------ | ----------------- | ------------ | ------------- | ------------------- | --------------------- | ------- | ------ |
83
| triton | standard | all<sup>1</sup> | G,A,T | silu, gelu,</br>swigluoai,</br>silu_no_mul,</br>gelu_no_mul | Y | Y | [`fused_experts`][vllm.model_executor.layers.fused_moe.fused_moe.fused_experts],</br>[`TritonExperts`][vllm.model_executor.layers.fused_moe.experts.triton_moe.TritonExperts] |
84
-| triton (batched) | batched | all<sup>1</sup> | G,A,T | silu, gelu | <sup>6</sup> | Y | [`BatchedTritonExperts`][vllm.model_executor.layers.fused_moe.fused_batched_moe.BatchedTritonExperts] |
84
+| triton (batched) | batched | all<sup>1</sup> | G,A,T | silu, gelu | <sup>6</sup> | Y | [`BatchedTritonExperts`][vllm.model_executor.layers.fused_moe.experts.fused_batched_moe.BatchedTritonExperts] |
85
| deep gemm | standard,</br>batched | fp8 | G(128),A,T | silu, gelu | <sup>6</sup> | Y | </br>[`DeepGemmExperts`][vllm.model_executor.layers.fused_moe.experts.deep_gemm_moe.DeepGemmExperts],</br>[`BatchedDeepGemmExperts`][vllm.model_executor.layers.fused_moe.experts.batched_deep_gemm_moe.BatchedDeepGemmExperts] |
86
| cutlass_fp4 | standard,</br>batched | nvfp4 | A,T | silu | Y | Y | [`CutlassExpertsFp4`][vllm.model_executor.layers.fused_moe.experts.cutlass_moe.CutlassExpertsFp4] |
87
| cutlass_fp8 | standard,</br>batched | fp8 | A,T | silu, gelu | Y | Y | [`CutlassExpertsFp8`][vllm.model_executor.layers.fused_moe.experts.cutlass_moe.CutlassExpertsFp8],</br>[`CutlasBatchedExpertsFp8`][vllm.model_executor.layers.fused_moe.experts.cutlass_moe.CutlassBatchedExpertsFp8] |
@@ -91,7 +91,7 @@ To be used with a particular `FusedMoEPrepareAndFinalizeModular` subclass, MoE k
91
| trtllm | standard | mxfp4,</br>nvfp4 | G(16),G(32) | <sup>5</sup> | N | Y | [`TrtLlmMxfp4ExpertsMonolithic`][vllm.model_executor.layers.fused_moe.experts.trtllm_mxfp4_moe.TrtLlmMxfp4ExpertsMonolithic],</br>[`TrtLlmMxfp4ExpertsModular`][vllm.model_executor.layers.fused_moe.experts.trtllm_mxfp4_moe.TrtLlmMxfp4ExpertsModular],</br>[`TrtLlmNvFp4ExpertsMonolithic`][vllm.model_executor.layers.fused_moe.experts.trtllm_nvfp4_moe.TrtLlmNvFp4ExpertsMonolithic],</br>[`TrtLlmNvfp4ExpertsModular`][vllm.model_executor.layers.fused_moe.experts.trtllm_nvfp4_moe.TrtLlmNvFp4ExpertsModular] |
92
| rocm aiter moe | standard | mxfp4,</br>fp8 | G(32),G(128),A,T | silu, gelu,</br>swigluoai | Y | N | `rocm_aiter_fused_experts`,</br>`AiterExperts` |
93
| cpu_fused_moe | standard | N/A | N/A | silu | N | N | [`CPUFusedMOE`][vllm.model_executor.layers.fused_moe.cpu_fused_moe.CPUFusedMOE] |
94
-| naive batched<sup>4</sup> | batched | int8,</br>fp8 | G,A,T | silu, gelu | <sup>6</sup> | Y | [`NaiveBatchedExperts`][vllm.model_executor.layers.fused_moe.fused_batched_moe.NaiveBatchedExperts] |
94
+| naive batched<sup>4</sup> | batched | int8,</br>fp8 | G,A,T | silu, gelu | <sup>6</sup> | Y | [`NaiveBatchedExperts`][vllm.model_executor.layers.fused_moe.experts.fused_batched_moe.NaiveBatchedExperts] |
95
96
!!! info "Table key"
97
1. All types: mxfp4, nvfp4, int4, int8, fp8
tests/kernels/moe/modular_kernel_tools/mk_objects.py
+4
-4
@@ -15,16 +15,16 @@ from vllm.model_executor.layers.fused_moe.experts.batched_deep_gemm_moe import (
15
BatchedDeepGemmExperts,
16
)
17
from vllm.model_executor.layers.fused_moe.experts.deep_gemm_moe import DeepGemmExperts
18
-from vllm.model_executor.layers.fused_moe.fused_batched_moe import (
18
+from vllm.model_executor.layers.fused_moe.experts.fused_batched_moe import (
19
BatchedTritonExperts,
20
NaiveBatchedExperts,
21
)
22
+from vllm.model_executor.layers.fused_moe.experts.triton_deep_gemm_moe import (
23
+ TritonOrDeepGemmExperts,
24
+)
25
from vllm.model_executor.layers.fused_moe.prepare_finalize import (
26
MoEPrepareAndFinalizeNoDPEPModular,
27
)
25
-from vllm.model_executor.layers.fused_moe.triton_deep_gemm_moe import (
26
- TritonOrDeepGemmExperts,
27
-)
28
from vllm.model_executor.layers.quantization.utils.nvfp4_utils import (
29
cutlass_fp4_supported,
30
)
tests/kernels/moe/test_batched_deepgemm.py
+1
-1
@@ -9,7 +9,7 @@ from vllm.model_executor.layers.fused_moe.config import fp8_w8a8_moe_quant_confi
9
from vllm.model_executor.layers.fused_moe.experts.batched_deep_gemm_moe import (
10
BatchedDeepGemmExperts,
11
)
12
-from vllm.model_executor.layers.fused_moe.fused_batched_moe import (
12
+from vllm.model_executor.layers.fused_moe.experts.fused_batched_moe import (
13
BatchedTritonExperts,
14
)
15
from vllm.model_executor.layers.fused_moe.modular_kernel import FusedMoEKernel
tests/kernels/moe/test_batched_moe.py
+1
-1
@@ -16,7 +16,7 @@ from tests.kernels.quant_utils import native_batched_masked_quant_matmul
16
from tests.kernels.utils import torch_experts
17
from vllm.config import VllmConfig, set_current_vllm_config
18
from vllm.model_executor.layers.fused_moe import fused_topk
19
-from vllm.model_executor.layers.fused_moe.fused_batched_moe import (
19
+from vllm.model_executor.layers.fused_moe.experts.fused_batched_moe import (
20
invoke_moe_batched_triton_kernel,
21
)
22
from vllm.platforms import current_platform
tests/kernels/moe/test_block_fp8.py
+1
-1
@@ -31,7 +31,7 @@ from vllm.model_executor.layers.fused_moe.config import (
31
from vllm.model_executor.layers.fused_moe.experts.deep_gemm_moe import (
32
_valid_deep_gemm_shape,
33
)
34
-from vllm.model_executor.layers.fused_moe.triton_deep_gemm_moe import (
34
+from vllm.model_executor.layers.fused_moe.experts.triton_deep_gemm_moe import (
35
TritonOrDeepGemmExperts,
36
)
37
from vllm.platforms import current_platform
tests/kernels/moe/test_deepep_moe.py
+3
-1
@@ -19,7 +19,9 @@ from vllm.model_executor.layers.fused_moe.activation import MoEActivation
19
from vllm.model_executor.layers.fused_moe.config import (
20
FusedMoEQuantConfig,
21
)
22
-from vllm.model_executor.layers.fused_moe.fused_batched_moe import BatchedTritonExperts
22
+from vllm.model_executor.layers.fused_moe.experts.fused_batched_moe import (
23
+ BatchedTritonExperts,
24
+)
25
from vllm.model_executor.layers.fused_moe.modular_kernel import FusedMoEKernel
26
from vllm.model_executor.layers.quantization.utils.fp8_utils import (
27
per_token_group_quant_fp8,
tests/kernels/moe/test_deepgemm.py
+2
-2
@@ -25,10 +25,10 @@ from vllm.model_executor.layers.fused_moe.config import (
25
FusedMoEQuantDesc,
26
fp8_w8a8_moe_quant_config,
27
)
28
-from vllm.model_executor.layers.fused_moe.fused_moe import fused_experts
29
-from vllm.model_executor.layers.fused_moe.triton_deep_gemm_moe import (
28
+from vllm.model_executor.layers.fused_moe.experts.triton_deep_gemm_moe import (
29
TritonOrDeepGemmExperts,
30
)
31
+from vllm.model_executor.layers.fused_moe.fused_moe import fused_experts
32
from vllm.model_executor.layers.quantization.utils.fp8_utils import (
33
per_token_group_quant_fp8,
34
)
tests/kernels/moe/utils.py
+4
-4
@@ -17,13 +17,13 @@ from vllm.model_executor.layers.fused_moe.config import (
17
FusedMoEQuantConfig,
18
RoutingMethodType,
19
)
20
-from vllm.model_executor.layers.fused_moe.experts.triton_moe import (
21
- TritonExperts,
22
-)
23
-from vllm.model_executor.layers.fused_moe.fused_batched_moe import (
20
+from vllm.model_executor.layers.fused_moe.experts.fused_batched_moe import (
21
BatchedTritonExperts,
22
NaiveBatchedExperts,
23
)
24
+from vllm.model_executor.layers.fused_moe.experts.triton_moe import (
25
+ TritonExperts,
26
+)
27
from vllm.model_executor.layers.fused_moe.fused_moe import (
28
fused_experts,
29
)
vllm/lora/layers/fused_moe.py
+1
-1
@@ -10,10 +10,10 @@ from vllm.config.lora import LoRAConfig
10
from vllm.distributed.utils import divide
11
from vllm.lora.layers.base import BaseLayerWithLoRA
12
from vllm.model_executor.layers.fused_moe import FusedMoE
13
+from vllm.model_executor.layers.fused_moe.experts.lora_context import MoELoRAContext
14
from vllm.model_executor.layers.fused_moe.fused_moe_modular_method import (
15
FusedMoEModularMethod,
16
)
16
-from vllm.model_executor.layers.fused_moe.lora_context import MoELoRAContext
17
from vllm.model_executor.layers.fused_moe.modular_kernel import FusedMoEKernel
18
from vllm.model_executor.layers.fused_moe.prepare_finalize import (
19
MoEPrepareAndFinalizeNoDPEPModular,
vllm/model_executor/layers/fused_moe/__init__.py
+6
-6
@@ -85,9 +85,15 @@ if HAS_TRITON:
85
from vllm.model_executor.layers.fused_moe.experts.deep_gemm_moe import (
86
DeepGemmExperts,
87
)
88
+ from vllm.model_executor.layers.fused_moe.experts.fused_batched_moe import (
89
+ BatchedTritonExperts,
90
+ )
91
from vllm.model_executor.layers.fused_moe.experts.rocm_aiter_moe import (
92
AiterExperts,
93
)
94
+ from vllm.model_executor.layers.fused_moe.experts.triton_deep_gemm_moe import (
95
+ TritonOrDeepGemmExperts,
96
+ )
97
from vllm.model_executor.layers.fused_moe.experts.triton_moe import (
98
TritonExperts,
99
TritonWNA16Experts,
@@ -97,9 +103,6 @@ if HAS_TRITON:
103
XPUExpertsFp8,
104
XPUExpertsMXFp4,
105
)
100
- from vllm.model_executor.layers.fused_moe.fused_batched_moe import (
101
- BatchedTritonExperts,
102
- )
106
from vllm.model_executor.layers.fused_moe.fused_moe import (
107
fused_experts,
108
get_config_file_name,
@@ -110,9 +113,6 @@ if HAS_TRITON:
113
from vllm.model_executor.layers.fused_moe.router.grouped_topk_router import (
114
GroupedTopk,
115
)
113
- from vllm.model_executor.layers.fused_moe.triton_deep_gemm_moe import (
114
- TritonOrDeepGemmExperts,
115
- )
116
117
__all__ += [
118
"AiterExperts",
vllm/model_executor/layers/fused_moe/experts/fallback.py
renamed
vllm/model_executor/layers/fused_moe/experts/fused_batched_moe.py
renamed
vllm/model_executor/layers/fused_moe/experts/fused_humming_moe.py
renamed
vllm/model_executor/layers/fused_moe/experts/gpt_oss_triton_kernels_moe.py
+3
-1
@@ -14,7 +14,9 @@ from vllm.model_executor.layers.fused_moe.config import (
14
FusedMoEQuantConfig,
15
RoutingMethodType,
16
)
17
-from vllm.model_executor.layers.fused_moe.lora_experts_mixin import LoRAExpertsMixin
17
+from vllm.model_executor.layers.fused_moe.experts.lora_experts_mixin import (
18
+ LoRAExpertsMixin,
19
+)
20
from vllm.model_executor.layers.fused_moe.topk_weight_and_reduce import (
21
TopKWeightAndReduceNoOP,
22
)
vllm/model_executor/layers/fused_moe/experts/lora_context.py
renamed
vllm/model_executor/layers/fused_moe/experts/lora_experts_mixin.py
renamed
+1
-1
@@ -3,7 +3,7 @@
3
4
import torch
5
6
-from vllm.model_executor.layers.fused_moe.lora_context import MoELoRAContext
6
+from vllm.model_executor.layers.fused_moe.experts.lora_context import MoELoRAContext
7
8
9
class LoRAExpertsMixin:
vllm/model_executor/layers/fused_moe/experts/marlin_moe.py
+3
-1
@@ -17,7 +17,9 @@ from vllm.model_executor.layers.fused_moe.config import (
17
FusedMoEParallelConfig,
18
FusedMoEQuantConfig,
19
)
20
-from vllm.model_executor.layers.fused_moe.lora_experts_mixin import LoRAExpertsMixin
20
+from vllm.model_executor.layers.fused_moe.experts.lora_experts_mixin import (
21
+ LoRAExpertsMixin,
22
+)
23
from vllm.model_executor.layers.fused_moe.moe_align_block_size import (
24
batched_moe_align_block_size,
25
moe_align_block_size,
vllm/model_executor/layers/fused_moe/experts/triton_cutlass_moe.py
renamed
+1
-1
@@ -11,8 +11,8 @@ from vllm.model_executor.layers.fused_moe.config import (
11
FusedMoEQuantConfig,
12
)
13
from vllm.model_executor.layers.fused_moe.experts.cutlass_moe import CutlassExpertsFp8
14
+from vllm.model_executor.layers.fused_moe.experts.fallback import FallbackExperts
15
from vllm.model_executor.layers.fused_moe.experts.triton_moe import TritonExperts
15
-from vllm.model_executor.layers.fused_moe.fallback import FallbackExperts
16
from vllm.platforms import current_platform
17
18
vllm/model_executor/layers/fused_moe/experts/triton_deep_gemm_moe.py
renamed
+1
-1
@@ -14,8 +14,8 @@ from vllm.model_executor.layers.fused_moe.experts.deep_gemm_moe import (
14
_valid_deep_gemm,
15
_valid_deep_gemm_shape,
16
)
17
+from vllm.model_executor.layers.fused_moe.experts.fallback import FallbackExperts
18
from vllm.model_executor.layers.fused_moe.experts.triton_moe import TritonExperts
18
-from vllm.model_executor.layers.fused_moe.fallback import FallbackExperts
19
from vllm.utils.deep_gemm import (
20
is_deep_gemm_e8m0_used,
21
)
vllm/model_executor/layers/fused_moe/experts/triton_moe.py
+3
-3
@@ -12,15 +12,15 @@ from vllm.model_executor.layers.fused_moe.config import (
12
FusedMoEParallelConfig,
13
FusedMoEQuantConfig,
14
)
15
+from vllm.model_executor.layers.fused_moe.experts.lora_experts_mixin import (
16
+ LoRAExpertsMixin,
17
+)
18
from vllm.model_executor.layers.fused_moe.fused_moe import (
19
_prepare_expert_assignment,
20
invoke_fused_moe_triton_kernel,
21
invoke_fused_moe_wna16_triton_kernel,
22
try_get_optimal_moe_config,
23
)
21
-from vllm.model_executor.layers.fused_moe.lora_experts_mixin import (
22
- LoRAExpertsMixin,
23
-)
24
from vllm.model_executor.layers.fused_moe.moe_align_block_size import (
25
moe_align_block_size,
26
)
vllm/model_executor/layers/fused_moe/oracle/fp8.py
+3
-3
@@ -130,7 +130,7 @@ def backend_to_kernel_cls(
130
return [FlashInferExperts]
131
132
elif backend == Fp8MoeBackend.DEEPGEMM:
133
- from vllm.model_executor.layers.fused_moe.triton_deep_gemm_moe import (
133
+ from vllm.model_executor.layers.fused_moe.experts.triton_deep_gemm_moe import (
134
TritonOrDeepGemmExperts,
135
)
136
@@ -158,7 +158,7 @@ def backend_to_kernel_cls(
158
return [TritonExperts]
159
160
elif backend == Fp8MoeBackend.BATCHED_TRITON:
161
- from vllm.model_executor.layers.fused_moe.fused_batched_moe import (
161
+ from vllm.model_executor.layers.fused_moe.experts.fused_batched_moe import (
162
BatchedTritonExperts,
163
)
164
@@ -172,7 +172,7 @@ def backend_to_kernel_cls(
172
return [AiterExperts]
173
174
elif backend == Fp8MoeBackend.VLLM_CUTLASS:
175
- from vllm.model_executor.layers.fused_moe.triton_cutlass_moe import (
175
+ from vllm.model_executor.layers.fused_moe.experts.triton_cutlass_moe import (
176
TritonOrCutlassExperts,
177
)
178
vllm/model_executor/layers/fused_moe/oracle/mxfp4.py
+1
-1
@@ -147,7 +147,7 @@ def backend_to_kernel_cls(
147
return [UnfusedOAITritonExperts]
148
149
elif backend == Mxfp4MoeBackend.HUMMING:
150
- from vllm.model_executor.layers.fused_moe.fused_humming_moe import (
150
+ from vllm.model_executor.layers.fused_moe.experts.fused_humming_moe import (
151
BatchedHummingGroupedExperts,
152
HummingGroupedExperts,
153
HummingIndexedExperts,
vllm/model_executor/layers/fused_moe/oracle/unquantized.py
+1
-1
@@ -116,7 +116,7 @@ def backend_to_kernel_cls(
116
return TritonExperts
117
118
elif backend == UnquantizedMoeBackend.BATCHED_TRITON:
119
- from vllm.model_executor.layers.fused_moe.fused_batched_moe import (
119
+ from vllm.model_executor.layers.fused_moe.experts.fused_batched_moe import (
120
BatchedTritonExperts,
121
)
122
vllm/model_executor/layers/quantization/humming.py
+2
-2
@@ -58,7 +58,7 @@ try:
58
)
59
from humming.utils.weight import quantize_weight
60
61
- from vllm.model_executor.layers.fused_moe.fused_humming_moe import (
61
+ from vllm.model_executor.layers.fused_moe.experts.fused_humming_moe import (
62
BatchedHummingGroupedExperts,
63
HummingGroupedExperts,
64
HummingIndexedExperts,
@@ -182,7 +182,7 @@ def compressed_tensors_get_config(config: dict[str, Any], key: str):
182
183
184
class HummingConfig(QuantizationConfig):
185
- packed_modules_mapping = {}
185
+ packed_modules_mapping: dict[str, list[str]] = {}
186
187
def __init__(self, full_config: dict[str, Any] | None = None):
188
assert_humming_available()
vllm/model_executor/warmup/deep_gemm_warmup.py
+2
-2
@@ -13,10 +13,10 @@ import vllm.envs as envs
13
from vllm.distributed.parallel_state import get_dp_group, is_global_first_rank
14
from vllm.model_executor.layers.fused_moe.deep_gemm_utils import compute_aligned_M
15
from vllm.model_executor.layers.fused_moe.experts.deep_gemm_moe import DeepGemmExperts
16
-from vllm.model_executor.layers.fused_moe.layer import FusedMoE
17
-from vllm.model_executor.layers.fused_moe.triton_deep_gemm_moe import (
16
+from vllm.model_executor.layers.fused_moe.experts.triton_deep_gemm_moe import (
17
TritonOrDeepGemmExperts,
18
)
19
+from vllm.model_executor.layers.fused_moe.layer import FusedMoE
20
from vllm.model_executor.layers.linear import LinearBase
21
from vllm.model_executor.layers.quantization.fp8 import Fp8LinearMethod
22
from vllm.model_executor.layers.quantization.online.mxfp8 import Mxfp8OnlineLinearMethod