From e3a85ec40d5a829b09aaa1778dd2700c5ee16b86 Mon Sep 17 00:00:00 2001 From: jmestwa-coder Date: Tue, 9 Jun 2026 10:33:07 +0530 Subject: [PATCH] bound remote device count in ib connect and accept Signed-off-by: jmestwa-coder --- src/transport/net_ib/connect.cc | 10 ++++++++++ 1 file changed, 10 insertions(+) diff --git a/src/transport/net_ib/connect.cc b/src/transport/net_ib/connect.cc index e2c8ecb35..3121e1d02 100644 --- a/src/transport/net_ib/connect.cc +++ b/src/transport/net_ib/connect.cc @@ -1025,6 +1025,11 @@ ncclResult_t ncclIbConnectImpl(void* ctx, int dev, void* opaqueHandle, void** se memcpy(&remMeta, stage->buffer, sizeof(ncclIbConnectionMetadata)); + if (remMeta.ndevs <= 0 || remMeta.ndevs > NCCL_IB_MAX_DEVS_PER_NIC) { + WARN("NET/IB : remote peer reported an invalid device count %d (max %d)", remMeta.ndevs, NCCL_IB_MAX_DEVS_PER_NIC); + return ncclInternalError; + } + // ensure that the remote devices have the same link layer than the local devices used in the connection. if (comm->base.vProps.ndevs > 0) { int ibDev0 = comm->devs[0].base.ibDevN; @@ -1459,6 +1464,11 @@ ncclResult_t ncclIbAcceptImpl(void* listenComm, void** recvComm, ncclNetDeviceHa /* copy back the received info */ memcpy(&remMeta, stage->buffer, sizeof(struct ncclIbConnectionMetadata)); + if (remMeta.ndevs <= 0 || remMeta.ndevs > NCCL_IB_MAX_DEVS_PER_NIC) { + WARN("NET/IB : remote peer reported an invalid device count %d (max %d)", remMeta.ndevs, NCCL_IB_MAX_DEVS_PER_NIC); + return ncclInternalError; + } + // Subnet-aware device selection: use the remote sender's GIDs to find a local // NIC on the same subnet. Override lComm->dev and update vProps if a // better device is found.