sglang-rocenante.patch 8.2 KB

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