From 243c064df2b5ce762d3eadddf8ec54ade796c6d8 Mon Sep 17 00:00:00 2001 From: Lianmin Zheng Date: Mon, 3 Nov 2025 16:11:00 -0800 Subject: [PATCH] Remove the dependency of nccl.h in symmetric memory (#12571) --- .../device_communicators/pynccl_allocator.py | 21 ++++++++++++++++++- 1 file changed, 20 insertions(+), 1 deletion(-) diff --git a/python/sglang/srt/distributed/device_communicators/pynccl_allocator.py b/python/sglang/srt/distributed/device_communicators/pynccl_allocator.py index 5a5e9ebe6..6f29185ca 100644 --- a/python/sglang/srt/distributed/device_communicators/pynccl_allocator.py +++ b/python/sglang/srt/distributed/device_communicators/pynccl_allocator.py @@ -8,10 +8,29 @@ from sglang.srt.distributed.parallel_state import GroupCoordinator from sglang.srt.server_args import get_global_server_args nccl_allocator_source = """ -#include + +#include extern "C" { +// copy from https://github.com/NVIDIA/nccl/blob/master/src/nccl.h.in +typedef enum { ncclSuccess = 0, + ncclUnhandledCudaError = 1, + ncclSystemError = 2, + ncclInternalError = 3, + ncclInvalidArgument = 4, + ncclInvalidUsage = 5, + ncclRemoteError = 6, + ncclInProgress = 7, + ncclNumResults = 8 } ncclResult_t; +typedef struct ncclComm* ncclComm_t; +typedef struct ncclWindow_vidmem* ncclWindow_t; +ncclResult_t ncclCommWindowRegister(ncclComm_t comm, void* buff, size_t size, ncclWindow_t* win, int winFlags); +#define NCCL_WIN_COLL_SYMMETRIC 0x01 + +ncclResult_t ncclMemAlloc(void** ptr, size_t size); +ncclResult_t ncclMemFree(void *ptr); + void* nccl_alloc_plug(size_t size, int device, void* stream) { void* ptr; ncclResult_t err = ncclMemAlloc(&ptr, size);