From 783c073a03db680aa16e88a83eb77e0b08db4f29 Mon Sep 17 00:00:00 2001 From: Sarat Kamisetty Date: Fri, 25 Jul 2025 20:56:58 +0530 Subject: [PATCH] passing down NET_OPTIONAL_RECV_COMPLETION hint to n/w plugin to enable optimizations (#1752) Co-authored-by: Sarat Kamisetty --- src/transport/net.cc | 10 ++++++++-- 1 file changed, 8 insertions(+), 2 deletions(-) diff --git a/src/transport/net.cc b/src/transport/net.cc index 4c858b299b..1669ac0679 100644 --- a/src/transport/net.cc +++ b/src/transport/net.cc @@ -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(