AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

passing down NET_OPTIONAL_RECV_COMPLETION hint to n/w plugin to enable optimizations (#1752)

Browse Source 
Co-authored-by: Sarat Kamisetty <sakamiset@amd.com>

[ROCm/rccl commit: 783c073a03]
This commit is contained in:
Sarat Kamisetty
2025-07-25 20:56:58 +05:30
committed by GitHub
parent b3a0cc5e96
commit 1719aa67be
1 changed files with 8 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rccl/src/transport/net.cc
+8 -2
View File
@@ -1360,8 +1360,14 @@ static ncclResult_t sendProxyProgress(struct ncclProxyState* proxyState, struct
// Coverity complains about the size here as pointing to an out-of-scope temporary. Which is nonsense,
// since size is a plain integer.
// coverity[use_invalid:FALSE]
NCCLCHECK(proxyState->ncclNet->isend(resources->netSendComm, buff, size, resources->tpRank, sub->sendMhandle, sub, sub->requests+buffSlot));
if (sub->requests[buffSlot] != NULL) {
void **requestPtr = sub->requests+buffSlot;
// for LL/LL128 protocols, completion event for write operation is not needed on the receiver side as
// the LL flags are actively polled to detect if full data is received or not, so this hint can be used
// by network plugin to optimize the transport for LL/LL128
bool ignoreCompletion = ncclParamNetOptionalRecvCompletion() && ((args->protocol == NCCL_PROTO_LL128) || (args->protocol == NCCL_PROTO_LL));
if (ignoreCompletion) *requestPtr = (void *)NCCL_NET_OPTIONAL_RECV_COMPLETION;
NCCLCHECK(proxyState->ncclNet->isend(resources->netSendComm, buff, size, resources->tpRank, sub->sendMhandle, sub, requestPtr));
if (*requestPtr != NULL) {
#if defined(ENABLE_NPKIT) && defined(ENABLE_NPKIT_EVENT_NET_SEND_ENTRY) && defined(ENABLE_NPKIT_EVENT_NET_SEND_EXIT)
NpKit::CollectCpuEvent(