Merge remote-tracking branch 'origin/v0.3.0-dev' into fix/iluvatar-triton-stub

Author: physics31415926

vllm_fl/__init__.py

@@ -1,5 +1,6 @@
 # Copyright (c) 2025 BAAI. All rights reserved.
 
+import importlib
 import os
 import logging
 import sys
@@ -77,12 +78,13 @@ def _patch_flash_attn_import():
 
 
 def _patch_custom_ops():
-    """Register torch.ops._C op schemas when vllm._C is unavailable."""
-    try:
-        import vllm._C  # noqa: F401
-        return
-    except (ImportError, OSError):
-        pass
+    """Register fallback schemas when neither vLLM extension ABI is present."""
+    for module_name in ("vllm._C", "vllm._C_stable_libtorch"):
+        try:
+            importlib.import_module(module_name)
+            return
+        except (ImportError, OSError):
+            continue
 
     try:
         import vllm_fl._C  # noqa: F401

vllm_fl/attention/utils.py

@@ -15,6 +15,14 @@ def patch_mm_encoder_attention():
     """
     import vllm.model_executor.layers.attention.mm_encoder_attention as mm_mod
     from vllm.v1.attention.backends.registry import AttentionBackendEnum
+    from vllm.platforms import current_platform
+
+    # PlatformFL is registered as an OOT platform, so CustomOp would otherwise
+    # bind MMEncoderAttention to forward_oot -> forward_native (Torch SDPA).
+    # NVIDIA must keep vLLM's CUDA dispatch so the selected FLASH_ATTN backend
+    # is actually used during multimodal profiling and inference.
+    if current_platform.is_cuda():
+        mm_mod.MMEncoderAttention.forward_oot = mm_mod.MMEncoderAttention.forward_cuda
 
     def _patched_maybe_get_vit_flash_attn_backend(attn_backend):
         if attn_backend == AttentionBackendEnum.FLASH_ATTN:

vllm_fl/dispatch/backends/vendor/metax/patches/pynccl_wrapper.py

@@ -463,7 +463,7 @@ def ncclCommWindowDeregister(self, comm: ncclComm_t, window: ncclWindow_t) -> No
 
 def _compile_friendly_all_reduce(self, input_):
     """Use torch.distributed all_reduce which Dynamo can trace."""
-    dist.all_reduce(input_, op=dist.ReduceOp.SUM)
+    dist.all_reduce(input_, op=dist.ReduceOp.SUM, group=self.device_group)
     return input_
 
 

vllm_fl/ops/fused_moe/fused_moe_utils.py

@@ -154,9 +154,13 @@ def _return_or_raise(
 
         return _return_or_raise(requested_backend, moe_config, activation_format)
 
-    # Handle explicit FlashInfer FP16 configuration.
-    if envs.is_set("VLLM_USE_FLASHINFER_MOE_FP16"):
-        if not envs.VLLM_USE_FLASHINFER_MOE_FP16:
+    # Handle explicit FlashInfer FP16 configuration. vLLM 0.24 removed
+    # this legacy name from vllm.envs, so read it directly from os.environ.
+    if "VLLM_USE_FLASHINFER_MOE_FP16" in os.environ:
+        use_flashinfer_moe_fp16 = os.environ[
+            "VLLM_USE_FLASHINFER_MOE_FP16"
+        ].strip().lower() in ("1", "true")
+        if not use_flashinfer_moe_fp16:
             if UnquantizedMoeBackend.FLASHINFER_TRTLLM in AVAILABLE_BACKENDS:
                 AVAILABLE_BACKENDS.remove(UnquantizedMoeBackend.FLASHINFER_TRTLLM)
             if UnquantizedMoeBackend.FLASHINFER_CUTLASS in AVAILABLE_BACKENDS:

vllm_fl/platform.py

@@ -10,12 +10,15 @@
 
 import torch
 
-# import custom ops, trigger op registration (CUDA only)
-try:
-    import vllm._C  # noqa
-    import vllm._C_stable_libtorch  # noqa
-except (ImportError, OSError):
-    pass  # NPU or other platforms may not have vllm._C
+# Import custom ops and trigger registration on both legacy and stable-ABI
+# vLLM wheels. Keep the attempts independent: official vLLM 0.24 wheels have
+# only the stable-ABI module, so failure of the legacy import must not skip it.
+import importlib
+for _extension in ("vllm._C", "vllm._C_stable_libtorch"):
+    try:
+        importlib.import_module(_extension)
+    except (ImportError, OSError):
+        pass  # Non-CUDA platforms may not ship either extension.
 
 from vllm.logger import init_logger
 from vllm.platforms import Platform, PlatformEnum
@@ -220,7 +223,8 @@ def check_and_update_config(cls, vllm_config: "VllmConfig") -> None:
             compilation_config.cudagraph_mode = CUDAGraphMode.PIECEWISE
 
         if (
-            parallel_config.data_parallel_size > 1
+            parallel_config.all2all_backend == "deepep_high_throughput"
+            and parallel_config.data_parallel_size > 1
             and compilation_config.cudagraph_mode != CUDAGraphMode.NONE
         ):
             # TODO: Piecewise Cuda graph might be enabled