⚠ 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