0 Related work
在diff from 1 to 7内，ceAlltoall使用了对称内存来完成数据的传输。需要参考这里的同步机制来实现VCCL机内不hang。
1 具体实现
1.1 nccl内可参考部分
diff from 1 to 7 2025-10-13 19.50.43.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
batchOpsParams
dst
void** dsts; void** srcs; size_t* sizes; size_t numOps; bool intraBatchSync;
ncclCeInitBatchOpsParams(…)
每个变量calloc成nRanks的平方
dstRank == comm→rank
dstRank != comm→rank
srcPtr = mySendBuff + dstRank * chunkBytes
dstPtr = myRecvBuff + selfRank→rank * chunkBytes
batchOpsParams
dst
dsts[batchOpsParams.numOps]=dstPtr srcs[batchOpsParams.numOps]=srcPtr sizes[batchOpsParams.numOps]=chunkBytes numOps++
chunkBytes = args→nElts * args→eltSize
for r+1 in range (nRanks):
batchOpsParams
dst
dsts[batchOpsParams.numOps]=peerRecvBuff srcs[batchOpsParams.numOps]=srcPtr sizes[batchOpsParams.numOps]=chunkBytes numOps++
peerRecvBuff = comm→devrState→lsaFlatBase + comm→devrState→bigSize * dstRank + args→recvWin→bigOffset + args→recvWin→userPtr
cudaMemcpyBatchAsync
batchParams
dst
operation
waitValue
operation address value/value64 flags alias
writeValue
operation address value/value64 flags alias
flushRemoteWrites
memoryBarrier
Data
ncclMemOpSync
nvls??
MCSync
UCSync
rank0
传递和等待的值是ceSeqNum++
multicast object
Host
Device
copy(H2D)
rank0
rank1
copy(H2D)
rank0
rank1
mcDstPtr
t1
t2
t2
rank1
rank0
rank1
mcDstPtr
t1
t2
t2
rank0
Host
Device
rank0
rank1
rank1
waitValue
waitValue
cuStreamBatchMemOp
stream
这里是cuStreamWaitValue32的batch版 一个rank可以一下放多个wait op
cuStreamBatchMemOp
stream
这里是cuStreamWaitValue32的batch版 一个rank可以一下放多个wait op
copyAsync
copyAsync
NCCL
change to complete
change to complete
cuStreamBatchMemOp
cuStreamBatchMemOp
rank0
Host
Device
rank0
same as above
rank1
rank2
rank0
t1
t2
t2时刻: symTeamObtain
Link to original
1.2 vccl方案
1.2.1 Use symmetric
Symmetric Memory P2p Design 2025-10-15 19.19.33.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
Host
Device
rank0
t1
rank0
rank1
cudaMemcpyAsync
copyH2D
Host
Device
rank1
rank1
waitValue op
wait
rank1 stream
rank0 stream
Host
Device
rank0
t2
copyH2D
rank0
rank1
Host
Device
rank1
rank1
waitValue op
use symmetric mem
Link to original
2 排期
准备cuda13.0开发环境和测试环境一天，开发大约1-2天，debug2天。总时间4-5天。