2d2c660104 · tduncan/Nvidia-Optimization-Stack-VLLM

tests/kernels/moe/modular_kernel_tools/common.py

-1

@@ -647,7 +647,6 @@ def make_modular_kernel(

647 modular_kernel = mk.FusedMoEKernel(

648 prepare_finalize=prepare_finalize,

649 fused_experts=fused_experts,

650 - inplace=False,

650 )

651

652 return modular_kernel

tests/kernels/moe/test_batched_deepgemm.py

-2

@@ -80,7 +80,6 @@ def test_batched_deepgemm_vs_triton(

80 mk_triton = FusedMoEKernel(

81 prep_finalize,

82 triton_experts,

83 - inplace=False,

83 )

84

85 out_triton = mk_triton.apply(

@@ -105,7 +104,6 @@ def test_batched_deepgemm_vs_triton(

104 mk_deepgemm = FusedMoEKernel(

105 prep_finalize,

106 deepgemm_experts,

108 - inplace=False,

107 )

108

109 out_deepgemm = mk_deepgemm.apply(

tests/kernels/moe/test_block_fp8.py

-1

@@ -269,7 +269,6 @@ def test_w8a8_block_fp8_deep_gemm_fused_moe(M, N, K, E, topk, seed, monkeypatch)

269 moe_config=moe_config,

270 quant_config=quant_config,

271 ),

272 - inplace=False,

272 )

273

274 def deep_gemm_moe_fp8(a, w1, w2, w1_s, w2_s, topk_weights, topk_ids):

tests/kernels/moe/test_cutlass_moe.py

-2

@@ -214,7 +214,6 @@ def run_with_expert_maps(

214 moe_config=moe_config,

215 quant_config=new_quant_config,

216 ),

217 - inplace=False,

217 )

218 out_tensor = out_tensor + kernel.apply(**kwargs)

219

@@ -284,7 +283,6 @@ def run_8_bit(

283 moe_config=moe_config,

284 quant_config=quant_config,

285 ),

287 - inplace=False,

286 )

287 return kernel.apply(**kwargs)

288

tests/kernels/moe/test_deepep_deepgemm_moe.py

-3

@@ -193,7 +193,6 @@ def make_ll_modular_kernel(

193 return FusedMoEKernel(

194 prepare_finalize=a2a,

195 fused_experts=fused_experts,

196 - inplace=False,

196 )

197

198

@@ -226,7 +225,6 @@ def make_ht_modular_kernel(

225 return FusedMoEKernel(

226 prepare_finalize=a2a,

227 fused_experts=fused_experts,

229 - inplace=False,

228 )

229

230

@@ -353,7 +351,6 @@ def triton_impl(

351 w2=w2,

352 topk_weights=topk_weights,

353 topk_ids=topk_ids,

356 - inplace=False,

354 quant_config=quant_config,

355 )

356

tests/kernels/moe/test_deepep_moe.py

-1

@@ -185,7 +185,6 @@ def make_modular_kernel(

185 mk = FusedMoEKernel(

186 prepare_finalize=a2a,

187 fused_experts=fused_experts,

188 - inplace=False,

188 )

189 return mk

190

tests/kernels/moe/test_deepgemm.py

-3

@@ -126,7 +126,6 @@ def run_single_case(m, n, k, topk, num_experts, block_size):

126 moe_config=moe_config,

127 quant_config=quant_config,

128 ),

129 - inplace=False,

129 )

130

131 # triton reference

@@ -136,7 +135,6 @@ def run_single_case(m, n, k, topk, num_experts, block_size):

135 w2=w2,

136 topk_weights=topk_weights,

137 topk_ids=topk_ids,

139 - inplace=False,

138 quant_config=quant_config,

139 )

140

@@ -323,7 +321,6 @@ def run_single_fp4_case(m, n, k, topk, num_experts):

321 moe_config=moe_config,

322 quant_config=quant_config,

323 ),

326 - inplace=False,

324 )

325

326 # DeepGEMM FP4 path

tests/kernels/moe/test_flashinfer.py

-3

@@ -233,7 +233,6 @@ def test_flashinfer_per_tensor_moe_fp8_no_graph(

233 td.w2_quantized,

234 topk_weights=topk_weights,

235 topk_ids=topk_ids,

236 - inplace=False,

236 activation=activation,

237 global_num_experts=e,

238 expert_map=None,

@@ -321,7 +320,6 @@ def test_flashinfer_cutlass_moe_fp8_no_graph(

320 td.w2_quantized,

321 topk_weights=topk_weights,

322 topk_ids=topk_ids,

324 - inplace=False,

323 activation=activation,

324 global_num_experts=e,

325 expert_map=None,

@@ -364,7 +362,6 @@ def test_flashinfer_cutlass_moe_fp8_no_graph(

362 moe_config=moe_config,

363 quant_config=quant_config,

364 ),

367 - inplace=False,

365 )

366

367 flashinfer_cutlass_output = kernel.apply(

tests/kernels/moe/test_flashinfer_b12x_moe.py

-1

@@ -200,7 +200,6 @@ def test_flashinfer_b12x_moe(

200 use_monolithic=False,

201 ),

202 experts,

203 - inplace=False,

203 )

204

205 score = torch.randn((m, e), device="cuda", dtype=dtype)

tests/kernels/moe/test_flashinfer_moe.py

-1

@@ -117,7 +117,6 @@ def test_flashinfer_fp4_moe_no_graph(

117 use_monolithic=False,

118 ),

119 FlashInferExperts(moe_config=moe_config, quant_config=quant_config),

120 - inplace=False,

120 )

121

122 flashinfer_output = flashinfer_experts.apply(

tests/kernels/moe/test_modular_oai_triton_moe.py

-1

@@ -190,7 +190,6 @@ def oai_triton_moe_impl(

190 use_monolithic=False,

191 ),

192 fused_experts,

193 - inplace=False,

193 )

194

195 return mk.apply(

tests/kernels/moe/test_moe_layer.py

-1

@@ -1159,7 +1159,6 @@ def make_fake_moe_layer(

1159 quant_config=quant_config,

1160 topk_weights=topk_weights,

1161 topk_ids=topk_ids,

1162 - inplace=False,

1162 activation=activation,

1163 apply_router_weight_on_input=apply_router_weight_on_input,

1164 global_num_experts=global_num_experts,

tests/kernels/moe/test_nvfp4_moe.py

-2

@@ -105,7 +105,6 @@ def test_cutlass_fp4_moe_no_graph(

105 moe_config=moe_config,

106 quant_config=quant_config,

107 ),

108 - inplace=False,

108 )

109

110 cutlass_output = kernel.apply(

@@ -226,7 +225,6 @@ def test_cutlass_fp4_moe_swiglustep(

225 moe_config=make_dummy_moe_config(),

226 quant_config=quant_config,

227 ),

229 - inplace=False,

228 )

229

230 cutlass_output = kernel.apply(

tests/kernels/moe/test_trtllm_nvfp4_moe.py

-1

@@ -201,7 +201,6 @@ def test_trtllm_fp4_moe_no_graph(

201 use_monolithic=False,

202 ),

203 trtllm_inner,

204 - inplace=False,

204 )

205

206 trtllm_output = trtllm_experts.apply(

tests/kernels/moe/utils.py

-3

@@ -142,7 +142,6 @@ def batched_moe(

142 quant_config=quant_config,

143 moe_config=moe_config,

144 ),

145 - inplace=False,

145 )

146

147 return fused_experts.apply(

@@ -195,7 +194,6 @@ def naive_batched_moe(

194 quant_config=quant_config,

195 moe_config=moe_config,

196 ),

198 - inplace=False,

197 )

198

199 return fused_experts.apply(

@@ -631,7 +629,6 @@ def modular_triton_fused_moe(

629 use_monolithic=False,

630 ),

631 TritonExperts(moe_config, quant_config),

634 - inplace=False,

632 )

633

634

vllm/model_executor/layers/fused_moe/config.py

-6

@@ -1296,12 +1296,6 @@ class FusedMoEConfig:

1296 # cannot silently select one and drop the clamp.

1297 swiglu_limit: float | None = None

1298

1299 - # This flag is used to disable the inplace optimization

1300 - # in MoE kernels. If this flag is True then the kernel

1301 - # should not be using inplace. If the flag is false, the

1302 - # kernel is free to use inplace or not.

1303 - disable_inplace: bool = True

1304 -

1299 def __post_init__(self):

1300 if self.dp_size > 1:

1301 logger.debug_once(

vllm/model_executor/layers/fused_moe/eep_reconfigure.py

-1

@@ -118,6 +118,5 @@ def make_eep_staged_quant_method(

118 mk.FusedMoEKernel(

119 prepare_finalize,

120 experts,

121 - inplace=moe_kernel.inplace,

121 ),

122 )

vllm/model_executor/layers/fused_moe/experts/marlin_moe.py

+1 -13

@@ -30,7 +30,6 @@ from vllm.model_executor.layers.fused_moe.topk_weight_and_reduce import (

30 )

31 from vllm.model_executor.layers.fused_moe.utils import (

32 _resize_cache,

33 - disable_inplace,

33 swiglu_limit_func,

34 )

35 from vllm.model_executor.layers.quantization.utils.marlin_utils import (

@@ -257,7 +256,6 @@ def fused_marlin_moe(

256 is_k_full: bool = True,

257 output: torch.Tensor | None = None,

258 input_dtype: torch.dtype | None = None,

260 - inplace: bool = False,

259 clamp_limit: float | None = None,

260 ) -> torch.Tensor:

261 """

@@ -286,10 +284,6 @@ def fused_marlin_moe(

284 - torch.Tensor: The output tensor after applying the MoE layer.

285 """

286

289 - if inplace:

290 - assert output is None, "Conflicting request"

291 - assert not disable_inplace()

292 -

287 quant_type = ScalarType.from_id(quant_type_id)

288 assert quant_type in [

289 scalar_types.uint4,

@@ -379,7 +373,7 @@ def fused_marlin_moe(

373 ).view(-1, topk, K)

374

375 if output is None:

382 - output = hidden_states if inplace else torch.empty_like(hidden_states)

376 + output = torch.empty_like(hidden_states)

377

378 if moe_sum is None:

379 return torch.sum(moe_output.view(-1, topk, K), dim=1, out=output)

@@ -417,7 +411,6 @@ def batched_fused_marlin_moe(

411 is_k_full: bool = True,

412 output: torch.Tensor | None = None,

413 input_dtype: torch.dtype | None = None,

420 - inplace: bool = False,

414 clamp_limit: float | None = None,

415 ) -> torch.Tensor:

416 """

@@ -449,8 +442,6 @@ def batched_fused_marlin_moe(

442 f"hidden states must be batched. e.g. [B, MAX_TOKENS, K]."

443 f"But got {hidden_states.size()}"

444 )

452 - if inplace:

453 - assert output is None, "Conflicting request."

445

446 quant_type = ScalarType.from_id(quant_type_id)

447 assert quant_type in [

@@ -508,9 +499,6 @@ def batched_fused_marlin_moe(

499 expert_num_tokens=expert_num_tokens,

500 )

501

511 - if output is None and inplace:

512 - output = hidden_states

513 -

502 # TODO (varun): This can be avoided by plumbing the marlin kernel to

503 # ignore topk_weights when topk_weights_ptr is a nullptr.

504 topk_weights = torch.ones(

vllm/model_executor/layers/fused_moe/experts/rocm_aiter_moe.py

-2

@@ -17,7 +17,6 @@ from vllm.model_executor.layers.fused_moe.config import (

17 from vllm.model_executor.layers.fused_moe.topk_weight_and_reduce import (

18 TopKWeightAndReduceNoOP,

19 )

20 -from vllm.model_executor.layers.fused_moe.utils import disable_inplace

20 from vllm.model_executor.layers.quantization.utils.quant_utils import (

21 QuantKey,

22 kFp8Dynamic128Sym,

@@ -516,7 +515,6 @@ class AiterExperts(mk.FusedMoEExpertsModular):

515 and output.is_contiguous()

516 and result.is_contiguous()

517 and output._base is None

519 - and disable_inplace()

518 ):

519 output.set_(result)

520 else:

vllm/model_executor/layers/fused_moe/fused_moe.py

+7 -124

@@ -5,7 +5,6 @@

5 import functools

6 import json

7 import os

8 -from collections.abc import Callable

8 from typing import Any

9

10 import torch

@@ -27,7 +26,6 @@ from vllm.model_executor.layers.fused_moe.moe_align_block_size import (

26 moe_align_block_size,

27 )

28 from vllm.model_executor.layers.fused_moe.utils import (

30 - disable_inplace,

29 moe_kernel_quantize_input,

30 )

31 from vllm.platforms import current_platform

@@ -1334,99 +1332,7 @@ def try_get_optimal_moe_config(

1332 return config

1333

1334

1337 -def inplace_fused_experts(

1338 - hidden_states: torch.Tensor,

1339 - w1: torch.Tensor,

1340 - w2: torch.Tensor,

1341 - topk_weights: torch.Tensor,

1342 - topk_ids: torch.Tensor,

1343 - activation: str = "silu",

1344 - apply_router_weight_on_input: bool = False,

1345 - use_fp8_w8a8: bool = False,

1346 - use_int8_w8a8: bool = False,

1347 - use_int8_w8a16: bool = False,

1348 - use_int4_w4a16: bool = False,

1349 - ocp_mx_scheme: str | None = None,

1350 - per_channel_quant: bool = False,

1351 - global_num_experts: int = -1,

1352 - expert_map: torch.Tensor | None = None,

1353 - w1_scale: torch.Tensor | None = None,

1354 - w2_scale: torch.Tensor | None = None,

1355 - w1_zp: torch.Tensor | None = None,

1356 - w2_zp: torch.Tensor | None = None,

1357 - a1_scale: torch.Tensor | None = None,

1358 - a2_scale: torch.Tensor | None = None,

1359 - block_shape: list[int] | None = None,

1360 - w1_bias: torch.Tensor | None = None,

1361 - w2_bias: torch.Tensor | None = None,

1362 -) -> None:

1363 - fused_experts_impl(

1364 - hidden_states,

1365 - w1,

1366 - w2,

1367 - topk_weights,

1368 - topk_ids,

1369 - True,

1370 - activation,

1371 - apply_router_weight_on_input,

1372 - use_fp8_w8a8,

1373 - use_int8_w8a8,

1374 - use_int8_w8a16,

1375 - use_int4_w4a16,

1376 - ocp_mx_scheme,

1377 - per_channel_quant,

1378 - global_num_experts,

1379 - expert_map,

1380 - w1_scale,

1381 - w2_scale,

1382 - w1_zp,

1383 - w2_zp,

1384 - a1_scale,

1385 - a2_scale,

1386 - block_shape,

1387 - w1_bias,

1388 - w2_bias,

1389 - )

1390 -

1391 -

1392 -def inplace_fused_experts_fake(

1393 - hidden_states: torch.Tensor,

1394 - w1: torch.Tensor,

1395 - w2: torch.Tensor,

1396 - topk_weights: torch.Tensor,

1397 - topk_ids: torch.Tensor,

1398 - activation: str = "silu",

1399 - apply_router_weight_on_input: bool = False,

1400 - use_fp8_w8a8: bool = False,

1401 - use_int8_w8a8: bool = False,

1402 - use_int8_w8a16: bool = False,

1403 - use_int4_w4a16: bool = False,

1404 - ocp_mx_scheme: str | None = None,

1405 - per_channel_quant: bool = False,

1406 - global_num_experts: int = -1,

1407 - expert_map: torch.Tensor | None = None,

1408 - w1_scale: torch.Tensor | None = None,

1409 - w2_scale: torch.Tensor | None = None,

1410 - w1_zp: torch.Tensor | None = None,

1411 - w2_zp: torch.Tensor | None = None,

1412 - a1_scale: torch.Tensor | None = None,

1413 - a2_scale: torch.Tensor | None = None,

1414 - block_shape: list[int] | None = None,

1415 - w1_bias: torch.Tensor | None = None,

1416 - w2_bias: torch.Tensor | None = None,

1417 -) -> None:

1418 - pass

1419 -

1420 -

1421 -direct_register_custom_op(

1422 - op_name="inplace_fused_experts",

1423 - op_func=inplace_fused_experts,

1424 - mutates_args=["hidden_states"],

1425 - fake_impl=inplace_fused_experts_fake,

1426 -)

1427 -

1428 -

1429 -def outplace_fused_experts(

1335 +def fused_experts_op(

1336 hidden_states: torch.Tensor,

1337 w1: torch.Tensor,

1338 w2: torch.Tensor,

@@ -1458,7 +1364,6 @@ def outplace_fused_experts(

1364 w2,

1365 topk_weights,

1366 topk_ids,

1461 - False,

1367 activation,

1368 apply_router_weight_on_input,

1369 use_fp8_w8a8,

@@ -1481,7 +1386,7 @@ def outplace_fused_experts(

1386 )

1387

1388

1484 -def outplace_fused_experts_fake(

1389 +def fused_experts_op_fake(

1390 hidden_states: torch.Tensor,

1391 w1: torch.Tensor,

1392 w2: torch.Tensor,

@@ -1511,28 +1416,12 @@ def outplace_fused_experts_fake(

1416

1417

1418 direct_register_custom_op(

1514 - op_name="outplace_fused_experts",

1515 - op_func=outplace_fused_experts,

1516 - fake_impl=outplace_fused_experts_fake,

1419 + op_name="fused_experts",

1420 + op_func=fused_experts_op,

1421 + fake_impl=fused_experts_op_fake,

1422 )

1423

1424

1520 -def torch_vllm_inplace_fused_experts(**kwargs) -> torch.Tensor:

1521 - torch.ops.vllm.inplace_fused_experts(**kwargs)

1522 - hidden_states = kwargs["hidden_states"]

1523 - return hidden_states

1524 -

1525 -

1526 -def torch_vllm_outplace_fused_experts(**kwargs) -> torch.Tensor:

1527 - return torch.ops.vllm.outplace_fused_experts(**kwargs)

1528 -

1529 -

1530 -def dispatch_fused_experts_func(inplace: bool) -> Callable[..., torch.Tensor]:

1531 - if inplace:

1532 - return torch_vllm_inplace_fused_experts

1533 - return torch_vllm_outplace_fused_experts

1534 -

1535 -

1425 def _prepare_expert_assignment(

1426 topk_ids: torch.Tensor,

1427 config: dict[str, Any],

@@ -1582,15 +1471,12 @@ def _prepare_expert_assignment(

1471 )

1472

1473

1585 -# TODO (bnell): replace this with modular op. Can get rid of inplace/outplace

1586 -# torch ops.

1474 def fused_experts(

1475 hidden_states: torch.Tensor,

1476 w1: torch.Tensor,

1477 w2: torch.Tensor,

1478 topk_weights: torch.Tensor,

1479 topk_ids: torch.Tensor,

1593 - inplace: bool = False,

1480 activation: MoEActivation = MoEActivation.SILU,

1481 apply_router_weight_on_input: bool = False,

1482 global_num_experts: int = -1,

@@ -1601,9 +1487,7 @@ def fused_experts(

1487 if quant_config is None:

1488 quant_config = FUSED_MOE_UNQUANTIZED_CONFIG

1489

1604 - assert not inplace or not disable_inplace()

1605 -

1606 - return dispatch_fused_experts_func(inplace)(

1490 + return torch.ops.vllm.fused_experts(

1491 hidden_states=hidden_states,

1492 w1=w1,

1493 w2=w2,

@@ -1667,7 +1551,6 @@ def fused_experts_impl(

1551 w2: torch.Tensor,

1552 topk_weights: torch.Tensor,

1553 topk_ids: torch.Tensor,

1670 - inplace: bool,

1554 activation: str = "silu",

1555 apply_router_weight_on_input: bool = False,

1556 use_fp8_w8a8: bool = False,

@@ -1775,7 +1658,7 @@ def fused_experts_impl(

1658 else:

1659 raise ValueError(f"Unsupported compute_type: {hidden_states.dtype}")

1660

1778 - out_hidden_states = hidden_states if inplace else torch.empty_like(hidden_states)

1661 + out_hidden_states = torch.empty_like(hidden_states)

1662

1663 qhidden_states, a1q_scale = moe_kernel_quantize_input(

1664 A=hidden_states,

vllm/model_executor/layers/fused_moe/fused_moe_modular_method.py

-2

@@ -51,14 +51,12 @@ class FusedMoEModularMethod(FusedMoEMethodBase, CustomOp):

51 moe_layer: torch.nn.Module,

52 old_quant_method: FusedMoEMethodBase,

53 prepare_finalize: FusedMoEPrepareAndFinalizeModular,

54 - inplace: bool = False,

54 ) -> "FusedMoEModularMethod":

55 return FusedMoEModularMethod(

56 old_quant_method,

57 FusedMoEKernel(

58 prepare_finalize,

59 old_quant_method.select_gemm_impl(prepare_finalize, moe_layer),

61 - inplace=inplace,

60 ),

61 )

62

vllm/model_executor/layers/fused_moe/layer.py

-5

@@ -50,9 +50,6 @@ from vllm.model_executor.layers.fused_moe.runner.shared_experts import (

50 from vllm.model_executor.layers.fused_moe.unquantized_fused_moe_method import (

51 UnquantizedFusedMoEMethod,

52 )

53 -from vllm.model_executor.layers.fused_moe.utils import (

54 - disable_inplace,

55 -)

53 from vllm.model_executor.layers.quantization.base_config import (

54 QuantizationConfig,

55 )

@@ -337,7 +334,6 @@ class FusedMoE(PluggableLayer):

334 routing_method=self.routing_method_type,

335 swiglu_limit=swiglu_limit,

336 # TODO: in_dtype == out_dtype?

340 - disable_inplace=disable_inplace() or shared_experts is not None,

337 )

338 if self.moe_config.use_mori_kernels:

339 assert self.rocm_aiter_fmoe_enabled, (

@@ -473,7 +469,6 @@ class FusedMoE(PluggableLayer):

469 self,

470 self.base_quant_method,

471 prepare_finalize,

476 - inplace=not self.moe_config.disable_inplace,

472 )

473 )

474

vllm/model_executor/layers/fused_moe/modular_kernel.py

+1 -19

@@ -27,7 +27,6 @@ from vllm.model_executor.layers.fused_moe.runner.shared_experts import (

27 )

28 from vllm.model_executor.layers.fused_moe.utils import (

29 _resize_cache,

30 - disable_inplace,

30 )

31 from vllm.model_executor.layers.quantization.utils.quant_utils import (

32 QuantKey,

@@ -1025,11 +1024,9 @@ class FusedMoEKernelModularImpl:

1024 self,

1025 prepare_finalize: FusedMoEPrepareAndFinalizeModular,

1026 fused_experts: FusedMoEExpertsModular,

1028 - inplace: bool = False,

1027 ):

1028 self.prepare_finalize = prepare_finalize

1029 self.fused_experts = fused_experts

1032 - self.inplace = inplace

1030 moe_parallel_config = fused_experts.moe_config.moe_parallel_config

1031 self.moe_parallel_config = moe_parallel_config

1032 self.is_dp_ep = (

@@ -1392,12 +1389,7 @@ class FusedMoEKernelModularImpl:

1389 Returns:

1390 - torch.Tensor: The output tensor after applying the MoE layer.

1391 """

1395 - if self.inplace:

1396 - assert shared_experts is None

1397 - assert not disable_inplace()

1398 - output = hidden_states

1399 - else:

1400 - output = torch.empty_like(hidden_states)

1392 + output = torch.empty_like(hidden_states)

1393

1394 local_num_experts = w1.shape[0]

1395 if global_num_experts == -1:

@@ -1473,7 +1465,6 @@ class FusedMoEKernelMonolithicImpl:

1465 that have fused router + experts (e.g. FLASHINFER_TRTLLM).

1466 """

1467

1476 - # TODO(rob): add inplace support.

1468 a1q, a1q_scale, router_logits = self.prepare_finalize.prepare(

1469 hidden_states,

1470 router_logits=router_logits,

@@ -1509,7 +1500,6 @@ class FusedMoEKernel:

1500 self,

1501 prepare_finalize: FusedMoEPrepareAndFinalize,

1502 fused_experts: FusedMoEExperts,

1512 - inplace: bool = False,

1503 ):

1504 super().__init__()

1505

@@ -1521,13 +1511,11 @@ class FusedMoEKernel:

1511 self.impl = FusedMoEKernelModularImpl(

1512 prepare_finalize,

1513 fused_experts,

1524 - inplace,

1514 )

1515

1516 elif isinstance(

1517 prepare_finalize, FusedMoEPrepareAndFinalizeMonolithic

1518 ) and isinstance(fused_experts, FusedMoEExpertsMonolithic):

1530 - assert not inplace

1519 self.impl = FusedMoEKernelMonolithicImpl(

1520 prepare_finalize,

1521 fused_experts,

@@ -1549,12 +1537,6 @@ class FusedMoEKernel:

1537 else:

1538 return False

1539

1552 - @property

1553 - def inplace(self) -> bool:

1554 - if isinstance(self.impl, FusedMoEKernelModularImpl):

1555 - return self.impl.inplace

1556 - return False

1557 -

1540 @property

1541 def is_monolithic(self) -> bool:

1542 return isinstance(self.impl, FusedMoEKernelMonolithicImpl)

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

-4

@@ -627,10 +627,6 @@ def make_fp8_moe_kernel(

627 kernel = mk.FusedMoEKernel(

628 prepare_finalize,

629 experts,

630 - inplace=(

631 - not moe_config.disable_inplace

632 - and fp8_backend != Fp8MoeBackend.FLASHINFER_CUTLASS

633 - ),

630 )

631

632 return kernel

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

-1

@@ -213,7 +213,6 @@ def make_int8_moe_kernel(

213 kernel = mk.FusedMoEKernel(

214 prepare_finalize,

215 experts,

216 - inplace=not moe_config.disable_inplace,

216 )

217

218 return kernel

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

-1

@@ -202,7 +202,6 @@ def make_wna16_moe_kernel(

202 return mk.FusedMoEKernel(

203 prepare_finalize,

204 experts,

205 - inplace=not moe_config.disable_inplace,

205 )

206

207

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

-3

@@ -1720,9 +1720,6 @@ def make_mxfp4_moe_kernel(

1720 kernel = mk.FusedMoEKernel(

1721 prepare_finalize,

1722 experts,

1723 - inplace=(

1724 - not moe_config.disable_inplace and mxfp4_backend not in TRTLLM_BACKENDS

1725 - ),

1723 )

1724

1725 return kernel

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

-2

@@ -555,8 +555,6 @@ def make_nvfp4_moe_kernel(

555 kernel = mk.FusedMoEKernel(

556 prepare_finalize,

557 experts,

558 - inplace=False,

558 )

559

561 - # TODO(rob): update inplace logic to be part of the kernel.

560 return kernel

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

-1

@@ -359,7 +359,6 @@ def make_unquantized_moe_kernel(

359 kernel = mk.FusedMoEKernel(

360 prepare_finalize,

361 experts,

362 - inplace=(not moe_config.disable_inplace and not is_monolithic),

362 )

363

364 return kernel

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

-1

@@ -192,5 +192,4 @@ def make_w4a8_moe_kernel(

192 return mk.FusedMoEKernel(

193 prepare_finalize,

194 experts,

195 - inplace=not moe_config.disable_inplace,

195 )

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

-1

@@ -355,7 +355,6 @@ def make_w4a8_int8_moe_kernel(

355 kernel = mk.FusedMoEKernel(

356 prepare_finalize,

357 experts,

358 - inplace=not moe_config.disable_inplace,

358 )

359

360 return kernel

vllm/model_executor/layers/fused_moe/prepare_finalize/deepep_ht.py

-2

@@ -383,7 +383,6 @@ class DeepEPHTPrepareAndFinalize(mk.FusedMoEPrepareAndFinalizeModular):

383 if event.event is not None:

384 event.current_stream_wait()

385 dbo_switch_to_comm()

386 - # Respect inplace outputs.

386 output.copy_(combined_x, non_blocking=True)

387

388 # TODO(lucas): refactor the modular kernel so this will be

@@ -394,7 +393,6 @@ class DeepEPHTPrepareAndFinalize(mk.FusedMoEPrepareAndFinalizeModular):

393 else:

394 # TODO(lucas): support this case with the refactored modular kernel

395 assert not dbo_enabled()

397 - # Respect inplace outputs.

396 output.copy_(combined_x, non_blocking=True)

397 return None

398

vllm/model_executor/layers/fused_moe/prepare_finalize/naive_dp_ep.py

-1

@@ -127,7 +127,6 @@ class MoEPrepareAndFinalizeNaiveDPEPModular(mk.FusedMoEPrepareAndFinalizeModular

127 assert topk == 1, (

128 "apply_router_weight_on_input is only implemented for topk=1"

129 )

130 - # Note: do not use inplace for shared experts overlap

130 a1 = a1 * topk_weights.to(a1.dtype)

131

132 a1q, scales = _quantize_and_setup_dispatch(a1, quant_config, defer_input_quant)

vllm/model_executor/layers/fused_moe/prepare_finalize/no_dp_ep.py

-1

@@ -71,7 +71,6 @@ class MoEPrepareAndFinalizeNoDPEPModular(mk.FusedMoEPrepareAndFinalizeModular):

71 assert topk == 1, (

72 "apply_router_weight_on_input is only implemented for topk=1"

73 )

74 - # Note: do not use inplace for shared experts overlap

74 a1 = a1 * topk_weights.to(a1.dtype)

75

76 a1q, a1q_scale = _quantize_input(a1, quant_config, defer_input_quant)

vllm/model_executor/layers/fused_moe/runner/shared_experts.py

-1

@@ -119,7 +119,6 @@ class SharedExperts:

119

120 if experts_order == SharedExpertsOrder.MULTI_STREAM_OVERLAPPED:

121 assert self._stream is not None

122 - assert self._moe_config.disable_inplace

122

123 # Record that the clone will be used by shared_experts_stream

124 # to avoid gc issue from deallocation of hidden_states_clone

vllm/model_executor/layers/fused_moe/utils.py

-10

@@ -1,6 +1,5 @@

1 # SPDX-License-Identifier: Apache-2.0

2 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project

3 -import functools

3 from math import prod

4

5 import torch

@@ -32,7 +31,6 @@ from vllm.model_executor.layers.quantization.utils.w8a8_utils import (

31 from vllm.platforms import current_platform

32 from vllm.triton_utils import tl, triton

33 from vllm.utils.math_utils import cdiv

35 -from vllm.utils.torch_utils import is_torch_equal_or_newer

34

35

36 @triton.jit

@@ -368,14 +366,6 @@ def normalize_batched_scales_shape(

366 return scales

367

368

371 -# Torch custom ops can't deal with outputs aliasing inputs so we need to

372 -# disable inplace for torch >= 2.9.

373 -# See https://github.com/vllm-project/vllm/issues/26378

374 -@functools.cache

375 -def disable_inplace() -> bool:

376 - return is_torch_equal_or_newer("2.9")

377 -

378 -

369 @triton.jit

370 def _pack_topk_ids_weights_kernel(

371 topk_ids_ptr,

vllm/model_executor/layers/quantization/bitsandbytes.py

-1

@@ -497,7 +497,6 @@ class BitsAndBytesMoEMethod(FusedMoEMethodBase):

497 w2=w2,

498 topk_weights=topk_weights,

499 topk_ids=topk_ids,

500 - inplace=not self.moe.disable_inplace,

500 activation=layer.activation,

501 apply_router_weight_on_input=layer.apply_router_weight_on_input,

502 global_num_experts=layer.global_num_experts,

vllm/model_executor/layers/quantization/compressed_tensors/compressed_tensors_moe/compressed_tensors_moe_wna16.py

-1

@@ -256,7 +256,6 @@ class CompressedTensorsWNA16MoEMethod(CompressedTensorsMoEMethod):

256 layer.w2_weight_packed,

257 topk_weights=topk_weights,

258 topk_ids=topk_ids,

259 - inplace=not self.moe.disable_inplace,

259 activation=layer.activation,

260 apply_router_weight_on_input=layer.apply_router_weight_on_input,

261 global_num_experts=layer.global_num_experts,

vllm/model_executor/layers/quantization/compressed_tensors/compressed_tensors_moe/compressed_tensors_moe_wna16_marlin.py

-1

@@ -573,5 +573,4 @@ class CompressedTensorsWNA16MarlinMoEMethod(CompressedTensorsMoEMethod):

573 workspace=layer.workspace,

574 input_dtype=self.marlin_input_dtype,

575 is_k_full=self.is_k_full,

576 - inplace=not self.moe.disable_inplace,

576 )

vllm/model_executor/layers/quantization/moe_wna16.py

-1

@@ -377,7 +377,6 @@ class MoeWNA16Method(FusedMoEMethodBase):

377 layer.w2_qweight,

378 topk_weights=topk_weights,

379 topk_ids=topk_ids,

380 - inplace=not self.moe.disable_inplace,

380 apply_router_weight_on_input=layer.apply_router_weight_on_input,

381 global_num_experts=layer.global_num_experts,

382 expert_map=layer.expert_map,

vllm/model_executor/layers/quantization/quark/quark_moe.py

-3

@@ -497,7 +497,6 @@ class QuarkW8A8Fp8MoEMethod(QuarkMoEMethod):

497 apply_router_weight_on_input=layer.apply_router_weight_on_input,

498 global_num_experts=layer.global_num_experts,

499 expert_map=layer.expert_map,

500 - inplace=not self.moe.disable_inplace,

500 )

501 else:

502 from vllm.model_executor.layers.fused_moe import fused_experts

@@ -508,7 +507,6 @@ class QuarkW8A8Fp8MoEMethod(QuarkMoEMethod):

507 w2=layer.w2_weight,

508 topk_weights=topk_weights,

509 topk_ids=topk_ids,

511 - inplace=not self.moe.disable_inplace,

510 activation=layer.activation,

511 apply_router_weight_on_input=layer.apply_router_weight_on_input,

512 global_num_experts=layer.global_num_experts,

@@ -809,7 +807,6 @@ class QuarkW8A8Int8MoEMethod(QuarkMoEMethod):

807 w2=layer.w2_weight,

808 topk_weights=topk_weights,

809 topk_ids=topk_ids,

812 - inplace=not self.moe.disable_inplace,

810 activation=layer.activation,

811 apply_router_weight_on_input=layer.apply_router_weight_on_input,

812 global_num_experts=layer.global_num_experts,

vllm/model_executor/models/arctic.py

-1

@@ -213,7 +213,6 @@ class ArcticMoE(nn.Module):

213 self.w2s,

214 topk_weights,

215 topk_ids,

216 - inplace=True,

216 )

217 if self.reduce_results and self.tp_size > 1:

218 final_hidden_states = tensor_model_parallel_all_reduce(final_hidden_states)

vllm/model_executor/models/bert_with_rope.py

+1 -1

@@ -326,7 +326,7 @@ class NomicMoE(nn.Module):

326 hidden_states, router_logits, self.top_k, renormalize=False

327 )

328

329 - final_hidden_states = torch.ops.vllm.outplace_fused_experts(

329 + final_hidden_states = torch.ops.vllm.fused_experts(

330 hidden_states=hidden_states,

331 w1=self.w1,

332 w2=self.w2,

vllm/model_executor/models/minicpm.py

+1 -1

@@ -181,7 +181,7 @@ class MiniCPMMoE(nn.Module):

181 )

182

183 final_hidden_states = fused_experts(

184 - hidden_states, self.ws, self.w2s, topk_weights, topk_ids, inplace=False

184 + hidden_states, self.ws, self.w2s, topk_weights, topk_ids

185 )

186

187 if self.tp_size > 1:

[MoE] Remove inplace fused experts mechanism (#43727)