Leoda

Support IntraNode P2p and Support Heterogeneous GPU Vendor in FlagCX

3 min read

0 architecture

FlagCX内对hostfunc有两种实现，都可用deviceAdaptor来正确调用，此外proxy progress被hostfunc正常触发之后实际的cudaMemcpyAsyc也有deviceAdaptor来正确调用。 但是与net.cc内不同的是机内p2p需要共享内存来完成数据的/同步信号的跨进程读写。综上整个p2p需要实现的话需要如下结构：

Support IntraNode P2p and Support Heterogeneous GPU Vendor in FlagCx 2025-10-20 15.32.45.excalidraw

⚠ Switch to EXCALIDRAW VIEW in the MORE OPTIONS menu of this document. ⚠ You can decompress Drawing data with the command palette: ‘Decompress current Excalidraw file’. For more info check in plugin settings under ‘Saving’

Excalidraw Data

Text Elements

deviceAdaptor→ hostfunc(semaphore)

group.cc

  1. task enqueue/dequeue 2. fill op message

hostfunc

event/stream + copy（H2D）+devicefunc

Device

Host

proxy progress

sender

receiver

semaphore→poll

semaphore→signalCounter

semaphore→poll

semaphore→signalCounter

cross process

same process

RecvFifo

sendBuff

recvBuff

HBM

sendrecv per op

sender/receiver

Shared Proxy Info

1.desc 2. SHM （sendHead recvTail ）

ncclShmAllocateShareableBuffer

ncclShmImportShareableBuffer

Shared Proxy Info

1.desc 2. SHM （sendHead recvTail ）

connectInfo (desc)

CPU Mem

bootstrapSend

setup

proxy

cuMemcpy

cuMemMap

cuMemCreate

cuMemAddressReserve

cuMemGetAllocationGranularity （granularity, prop）

cuMemSetAccess x 2 （CUmemAccessDesc.location.type）

virtual Address

handle

CUdeviceptr

desc

  1. setup+connect bootstapSend/Recv

ncclShmAllocateShareableBuffer

ncclShmImportShareableBuffer

cuMemImportFromShareableHandle fd变成handle

cuMemSetAccess x 2 （CUmemAccessDesc.location.type）

ncclProxyClientGetFdBlocking(…) 用desc表示handle的部分发给对端，对端export fd，再走UDS传回来

use cumem

use cudaIpc finish allocate/import shared buffer

Shared Proxy Info

RecvFifo

receiver(process2)

sender(process1)

cudaDeviceEnablePeerAccess （directPtr变成devMem）

p2pMap

p2pMap

cudaIpcGetMemHandle（*p2pBuff→directPtr）

psmP2pAllocateShareableBuffer

cudaMalloc(**p2pBuff→directPtr)

cudaIpcOpenMemHandle（devMemPtr）

psmP2pImportShareableBuffer

ncclShmAllocateShareableBuffer

ShmOpen

Shared Proxy Info

p2pShm

ConnectInfo.desc

fill

bootstrapSend p2pBuff/desc

ncclShmImportShareableBuffer

ShmOpen

reOpen

ConnectInfo.desc

dev/shm/

mmap

mmap

Shared Proxy Info

p2pShm

Link to original

1 FlagCX Adaptor

跨GPU通过bootstrap去交换的是 CUmemAllocationHandleType_enum，我们需要在flagcx内至少支持其中一类，比如 CU_MEM_HANDLE_TYPE_POSIX_FILE_DESCRIPTOR(fd) 。

此外，nvidia以外厂商需要支持如下类似虚拟内存管理的结构体和接口： 6.14 Virtual Memory Management(VMM)，调整为使用Ipc，不使用cuMem。

1.1 flagcx pr-272

flagcxDeviceAdaptor和对应实例内增加接口如下：

APInvidia
1ipcMemHandleOpencudaIpcOpenMemHandle
2ipcMemHandleClosecudaIpcCloseMemHandle
3ipcMemHandleGetcudaIpcGetMemHandle
4ipcMemHandleCreateflagcxCalloc
5ipcMemHandleFree直接free
sharedbuffer则是把nccl的shm.cc内的cuMemDisable情况的ncclShmAllocateShareableBuffer和ncclShmImportShareableBuffer搬了过来。这里实现的逻辑参考shm.cc。

1.2 还需要的

CUDA APIDescribe
1cudaThreadExchangeStreamCaptureMode放宽stream当前模式为宽松，允许cudaMalloc。防止当前在捕获期间，出现不安全cudaMalloc
2cudaDeviceEnablePeerAccessreceiver创建RecvFifo的时候需要设置当前GPU可被对端访问

2 Specific

2.1 分配GPU上的内存以及CPU上的物理页给两个进程使用

  1. 用proxy先去触发 让sender和receiver各自去准备自己需要共享的东西。
  2. 然后transport内connectInfo内填写desc，用bootstrap互相交换connectInfo。
  3. 再各自调用connect去把desc转换成各自需要的东西。

2.1.1 shm buffer

sender在自己的proxy setup阶段去调用shmutils.cc内的flagcxShmAllocateShareableBuffer。 ⚠️注意：

  1. 在flagcx内环境变量 FLAGCX_ENABLE_TOPO_DETECT 默认是关，就不会走initTransportsRank（fillPeerInfo，bootstrapAllGather等），所以就不会有 comm->peerInfo，会coredump。
Sender (rank 0):
  conn->proxyConn.connection->transportResources = resources_sender
    ├─ sendDevMem (自己的发送缓冲区，不用)
    ├─ recvDevMem (对方的接收缓冲区，导入的)
    └─ proxyInfo_sender
         ├─ shm (指向共享内存)
         ├─ desc（导receiver开辟的buffer）
         ├─ stream
         └─ events[MAXSTEPS--16]
 
Receiver (rank 1):
  conn->proxyConn.connection->transportResources = resources_receiver
    ├─ recvDevMem (自己的接收缓冲区，分配的)
    ├─ sendDevMem (对方的发送缓冲区，不用)
    └─ proxyInfo_receiver
         ├─ shm (指向共享内存)
         ├─ desc（自己的进程的不用导）
         ├─ stream
         └─ events[MAXSTEPS--16]

Support IntraNode P2p and Support Heterogeneous GPU Vendor in FlagCX 2025-10-30 18.44.29.excalidraw

⚠ Switch to EXCALIDRAW VIEW in the MORE OPTIONS menu of this document. ⚠ You can decompress Drawing data with the command palette: ‘Decompress current Excalidraw file’. For more info check in plugin settings under ‘Saving’

Excalidraw Data

Text Elements

Link to original

Graph View