rocenante.diff 9.3 KB

123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145146147148149150151152153154155156157158159160161162163164165166167168169170171172173174175176177178179180181182183184185186187188189190191192193194195196197198199200201202203204205206207208209210211212213214215216217218219220
  1. diff --git a/python/sglang/srt/distributed/device_communicators/pynccl.py b/python/sglang/srt/distributed/device_communicators/pynccl.py
  2. index 46aef5b945..fddacd7257 100644
  3. --- a/python/sglang/srt/distributed/device_communicators/pynccl.py
  4. +++ b/python/sglang/srt/distributed/device_communicators/pynccl.py
  5. @@ -3,7 +3,10 @@
  6. # Adapted from https://github.com/vllm-project/vllm/blob/v0.6.4.post1/vllm/distributed/device_communicators/pynccl.py
  7. import logging
  8. -from contextlib import contextmanager
  9. +import json
  10. +import os
  11. +import weakref
  12. +from contextlib import contextmanager, nullcontext
  13. from typing import Optional, Union
  14. # ===================== import region =====================
  15. @@ -27,6 +30,25 @@ from sglang.srt.utils.common import get_current_device_stream_fast
  16. logger = logging.getLogger(__name__)
  17. +_roce_runtimes = weakref.WeakSet()
  18. +
  19. +
  20. +def synchronize_and_check_roce(copy_done):
  21. + """Finish this result before checking transport health or consuming output."""
  22. + if copy_done is not None:
  23. + copy_done.synchronize()
  24. + elif _roce_runtimes:
  25. + # Non-overlap paths have no copy event. Complete all forward streams
  26. + # before reading the host failure word; an early check is insufficient.
  27. + torch.cuda.synchronize()
  28. + for runtime in tuple(_roce_runtimes):
  29. + runtime.check_health()
  30. + if not getattr(runtime, "_sglang_post_result_logged", False):
  31. + logger.info("ROCE_POST_RESULT_HEALTH %s", json.dumps(runtime.stats()))
  32. + runtime._sglang_post_result_logged = True
  33. +
  34. +
  35. +
  36. class PyNcclCommunicator:
  37. def __init__(
  38. self,
  39. @@ -34,6 +56,7 @@ class PyNcclCommunicator:
  40. device: Union[int, str, torch.device],
  41. library_path: Optional[str] = None,
  42. is_symmetric_memory_enabled: bool = False,
  43. + roce_enabled: bool = False,
  44. ):
  45. """
  46. Args:
  47. @@ -60,6 +83,8 @@ class PyNcclCommunicator:
  48. self.world_size = group.world_size
  49. self.group = group
  50. + self.roce = None
  51. + self._roce_route_logged = False
  52. # if world_size == 1, no need to create communicator
  53. if self.world_size == 1:
  54. @@ -131,11 +156,45 @@ class PyNcclCommunicator:
  55. warmup_stream.synchronize()
  56. del data
  57. + if roce_enabled and os.environ.get("SGLANG8_ROCE_ALLREDUCE", "0") == "1":
  58. + if not (2 <= self.world_size <= 16) or isinstance(group, StatelessProcessGroup):
  59. + raise RuntimeError("SG17 RoCEnante needs a Gloo exchange group of 2..16 ranks")
  60. + from b12x.comm import roce
  61. +
  62. + self.roce = roce.AllReduce.from_exchange_group(
  63. + exchange_group=group, device=self.device,
  64. + max_size=512 * 1024, max_gather_bytes=0,
  65. + )
  66. + self.roce.prepare((torch.bfloat16, torch.float32))
  67. + if len(self.roce.hca_names) != 2:
  68. + raise RuntimeError("SG17 RoCEnante requires two active RoCE interfaces")
  69. + _roce_runtimes.add(self.roce)
  70. + logger.info("ROCE_TP8_READY %s", json.dumps(self.roce.stats()))
  71. +
  72. +
  73. # by default it is disabled, e.g. in profiling models and prefill phase.
  74. # to use it, use under `with obj.change_state(enable=True)`, usually
  75. # when we are using CUDA graph.
  76. self.disabled = True
  77. + def roce_capture(self, stream):
  78. + return self.roce.capture(stream=stream) if self.roce is not None else nullcontext()
  79. +
  80. + def _should_roce(self, tensor, op):
  81. + if self.roce is None:
  82. + return False
  83. + self.roce.check_health()
  84. + eligible = (
  85. + op == ReduceOp.SUM
  86. + and tensor.dtype in (torch.bfloat16, torch.float32)
  87. + and self.roce.should_allreduce(tensor)
  88. + )
  89. + if eligible and not self._roce_route_logged:
  90. + logger.info("ROCE_TP8_ROUTE rank=%d dtype=%s bytes=%d", self.rank,
  91. + tensor.dtype, tensor.numel() * tensor.element_size())
  92. + self._roce_route_logged = True
  93. + return eligible
  94. +
  95. def _resolve_stream(self) -> torch.cuda.Stream:
  96. """Return the current device stream used for NCCL calls."""
  97. return get_current_device_stream_fast()
  98. @@ -151,6 +210,9 @@ class PyNcclCommunicator:
  99. f"but the input tensor is on {tensor.device}"
  100. )
  101. stream = self._resolve_stream()
  102. + if self._should_roce(tensor, op):
  103. + self.roce.all_reduce(tensor, out=tensor, stream=stream)
  104. + return
  105. self.nccl.ncclAllReduce(
  106. buffer_type(tensor.data_ptr()),
  107. buffer_type(tensor.data_ptr()),
  108. @@ -178,6 +240,8 @@ class PyNcclCommunicator:
  109. out_tensor = torch.empty_like(in_tensor)
  110. stream = self._resolve_stream()
  111. + if self._should_roce(in_tensor, op):
  112. + return self.roce.all_reduce(in_tensor, out=out_tensor, stream=stream)
  113. self.nccl.ncclAllReduce(
  114. buffer_type(in_tensor.data_ptr()), # sendbuff
  115. buffer_type(out_tensor.data_ptr()), # recvbuff - DIFFERENT pointer
  116. diff --git a/python/sglang/srt/distributed/parallel_state.py b/python/sglang/srt/distributed/parallel_state.py
  117. index e7bcc1d88e..efaef0215c 100644
  118. --- a/python/sglang/srt/distributed/parallel_state.py
  119. +++ b/python/sglang/srt/distributed/parallel_state.py
  120. @@ -468,6 +468,7 @@ class GroupCoordinator:
  121. group=self.cpu_group,
  122. device=self.device,
  123. is_symmetric_memory_enabled=self.is_symmetric_memory_enabled(),
  124. + roce_enabled=(group_name == "tp"),
  125. )
  126. self.pymscclpp_comm: Optional[PyMscclppCommunicator] = None
  127. @@ -654,7 +655,10 @@ class GroupCoordinator:
  128. maybe_pymscclpp_context = nullcontext()
  129. else:
  130. maybe_pymscclpp_context = pymscclpp_comm.change_state(enable=True)
  131. - with maybe_pynccl_context, maybe_pymscclpp_context:
  132. + maybe_roce_context = (
  133. + pynccl_comm.roce_capture(stream) if pynccl_comm else nullcontext()
  134. + )
  135. + with maybe_pynccl_context, maybe_pymscclpp_context, maybe_roce_context:
  136. yield graph_capture_context
  137. def all_reduce(self, input_: torch.Tensor) -> torch.Tensor:
  138. @@ -2059,6 +2063,8 @@ class GroupCoordinator:
  139. torch.distributed.destroy_process_group(self.cpu_group)
  140. self.cpu_group = None
  141. if self.pynccl_comm is not None:
  142. + if self.pynccl_comm.roce is not None:
  143. + self.pynccl_comm.roce.close()
  144. self.pynccl_comm = None
  145. if self.pymscclpp_comm is not None:
  146. self.pymscclpp_comm.destroy()
  147. diff --git a/python/sglang/srt/managers/scheduler_components/batch_result_processor.py b/python/sglang/srt/managers/scheduler_components/batch_result_processor.py
  148. index 7ce4572f59..86fc4e0508 100644
  149. --- a/python/sglang/srt/managers/scheduler_components/batch_result_processor.py
  150. +++ b/python/sglang/srt/managers/scheduler_components/batch_result_processor.py
  151. @@ -14,6 +14,8 @@ from typing import (
  152. import torch
  153. +from sglang.srt.distributed.device_communicators.pynccl import synchronize_and_check_roce
  154. +
  155. from sglang.srt.disaggregation.utils import DisaggregationMode
  156. from sglang.srt.environ import envs
  157. from sglang.srt.layers.logits_processor import (
  158. @@ -263,8 +265,7 @@ class SchedulerBatchResultProcessor:
  159. self.token_to_kv_pool_allocator.free_group_begin()
  160. if self.is_generation:
  161. - if result.copy_done is not None:
  162. - result.copy_done.synchronize()
  163. + synchronize_and_check_roce(result.copy_done)
  164. auxiliary_output_starts = self.snapshot_auxiliary_output_starts(
  165. batch, result
  166. )
  167. @@ -441,8 +442,7 @@ class SchedulerBatchResultProcessor:
  168. )
  169. else: # embedding or reward model
  170. - if result.copy_done is not None:
  171. - result.copy_done.synchronize()
  172. + synchronize_and_check_roce(result.copy_done)
  173. embeddings = self._convert_embeddings(result=result)
  174. phs = result.pooled_hidden_states
  175. @@ -868,8 +868,7 @@ class SchedulerBatchResultProcessor:
  176. is_decode = batch.forward_mode.is_decode()
  177. if not (is_decode or batch.forward_mode.is_extend()):
  178. return
  179. - if result.copy_done is not None:
  180. - result.copy_done.synchronize()
  181. + synchronize_and_check_roce(result.copy_done)
  182. next_token_ids = result.next_token_ids.tolist()
  183. if not is_decode:
  184. @@ -911,8 +910,7 @@ class SchedulerBatchResultProcessor:
  185. batch: ScheduleBatch,
  186. result: GenerationBatchResult,
  187. ):
  188. - if result.copy_done is not None:
  189. - result.copy_done.synchronize()
  190. + synchronize_and_check_roce(result.copy_done)
  191. self.output_streamer._stream_output_generation(
  192. batch.reqs, batch.return_logprob, is_idle_batch=True
  193. @@ -923,8 +921,7 @@ class SchedulerBatchResultProcessor:
  194. batch: ScheduleBatch,
  195. result: GenerationBatchResult,
  196. ):
  197. - if result.copy_done is not None:
  198. - result.copy_done.synchronize()
  199. + synchronize_and_check_roce(result.copy_done)
  200. auxiliary_output_starts = self.snapshot_auxiliary_output_starts(batch, result)
  201. auxiliary_output = result.auxiliary_host_output
  202. if result.routed_experts_output is not None: