本篇文章主要总结以下在传输路径方式选择的时候，选择每一种方式应该满足的条件和优先度。
本文初步总结，之后还会进行更新，欢迎大家补充
源码位置：tools/topo_expl
Topo结构：
初始化判断前 ret设置为0，代表此类方法不可使用，如果满足transports中的某一类，则设置 ret = 1.
PATH_LOC为节点到自己，PATH_NVL表示路径上的边都是NVLink，PATH_PIX表示经过最多一个PCIe switch，PATH_PXB表示经过了多个PCIe witch，但是没有经过CPU，PATH_PHB表示经过了CPU，PATH_SYS表示不同numa之间的路径。
# define
// Local (myself)
#define PATH_LOC 0
// Connection traversing NVLink
#define PATH_NVL 1
// Connection through NVLink using an intermediate GPU
#define PATH_NVB 2
// Connection traversing at most a single PCIe bridge
#define PATH_PIX 3
// Connection traversing multiple PCIe bridges (without traversing the PCIe Host Bridge)
#define PATH_PXB 4
// Connection between a GPU and a NIC using an intermediate GPU. Used to enable rail-local, aggregated network send/recv operations.
#define PATH_PXN 5
// Connection traversing PCIe as well as a PCIe Host Bridge (typically the CPU)
#define PATH_PHB 6
// Connection traversing PCIe as well as the SMP interconnect between NUMA nodes (e.g., QPI/UPI)
#define PATH_SYS 7
// Connection through the network
#define PATH_NET 8
// Disconnected
#define PATH_DIS 9P2P：
看一下源码的逻辑过程和源码
1. 检查设备是否支持精细粒度 GPU 操作
如果任一设备不支持 Fine-Grain 操作（精细粒度的 GPU 共享内存），直接判定两设备无法使用 P2P 通信。
2. 检查设备是否属于同一主机或容器
比较 hostHash 和 shmDev：
hostHash：表示设备所属主机的唯一标识。
shmDev：表示设备是否共享同一块内存设备（如共享同一块显存）。
如果设备不在同一主机或共享内存域，判定无法使用 P2P 通信。
3. 检查设备间的拓扑关系和通信路径
调用 ncclTopoCheckP2p 检查两设备之间的 P2P 拓扑关系：
如果 *ret = 0：表示设备间无直接 P2P 通信路径。
如果 intermediateRank != -1：表示设备间通信需要通过第三方设备中转。
如果 useMemcpy（标志位）为真，优先禁用 P2P 通信
4. 判断是否使用网络通信更优
调用 ncclTopoCheckNet 检查是否通过网络通信（如 NIC）会比 P2P 更优：
如果 useNet = 1：说明网络通信更适合，禁用 P2P。
/* Determine if two peers can communicate with P2P */
ncclResult_t p2pCanConnect(int* ret, struct ncclTopoSystem* topo, struct ncclTopoGraph* graph, struct ncclPeerInfo* info1, struct ncclPeerInfo* info2) {
if (!info1->hasFineGrain || !info2->hasFineGrain) {
*ret = 0;
return ncclSuccess;
}
// Rule out different nodes / isolated containers
if (info1->hostHash != info2->hostHash || info1->shmDev != info2->shmDev) {
*ret = 0;
return ncclSuccess;
}
// Check topology / p2p level.
int intermediateRank;
NCCLCHECK(ncclTopoCheckP2p(topo, info1->busId, info2->busId, ret, NULL, &intermediateRank));
if (*ret == 0) return ncclSuccess;
if (intermediateRank != -1) {
if (useMemcpy) *ret = 0;
return ncclSuccess;
}
// Check if NET would work better
int useNet = 0;
NCCLCHECK(ncclTopoCheckNet(topo, info1->busId, info2->busId, &useNet));
if (useNet) {
*ret = 0;
return ncclSuccess;
}
*ret = 1;
return ncclSuccess;
}
这里主要是再看一下ncclTopoCheckP2p，这个函数中会看Topo结构中两个GPU之间是否满足p2p，主要是p2p_Level 是否满足，看一下源码，如果CPU类型是 ARM CPUs/ CPU_ARCH_X86（Intel / ZHAOXIN）p2pLevel 有进一步的要求 p2pLevel = PATH_PXB，实际路径满足 path->type <= p2pLevel， 允许 p2p = 1.
源码位置是编译过后的位置：build/release/hipify/src/graph/paths.cc
int ncclTopoUserP2pLevel = -1;
ncclResult_t ncclTopoCheckP2p(struct ncclTopoSystem* system, int64_t id1, int64_t id2, int* p2p, int *read, int* intermediateRank) {
*p2p = 0;
if (read) *read = 0;
if (intermediateRank) *intermediateRank = -1;
// Get GPUs from topology
int g1, g2;
NCCLCHECK(ncclTopoIdToIndex(system, GPU, id1, &g1));
struct ncclTopoNode* gpu1 = system->nodes[GPU].nodes+g1;
if (ncclTopoIdToIndex(system, GPU, id2, &g2) == ncclInternalError) {
// GPU not found, we can't use p2p.
return ncclSuccess;
}
int intermediateIndex = -1;
// Set intermediate GPU rank, if routing through an intermediate GPU.
struct ncclTopoLinkList* path = gpu1->paths[GPU]+g2;
if (path->count == 2) {
struct ncclTopoNode* intermediateNode = path->list[0]->remNode;
if (intermediateNode->type == GPU) {
intermediateIndex = intermediateNode - system->nodes[GPU].nodes;
if (intermediateRank) *intermediateRank = intermediateNode->gpu.rank;
}
}
// In general, use P2P whenever we can.
int p2pLevel = PATH_SYS;
// User override
if (ncclTopoUserP2pLevel == -1)
NCCLCHECK(ncclGetLevel(&ncclTopoUserP2pLevel, "NCCL_P2P_DISABLE", "NCCL_P2P_LEVEL"));
if (ncclTopoUserP2pLevel != -2) {
p2pLevel = ncclTopoUserP2pLevel;
goto compare;
}
// Don't use P2P through ARM CPUs
int arch, vendor, model;
NCCLCHECK(ncclTopoCpuType(system, &arch, &vendor, &model));
if (arch == NCCL_TOPO_CPU_ARCH_ARM) p2pLevel = PATH_PXB;
if (arch == NCCL_TOPO_CPU_ARCH_X86 && vendor == NCCL_TOPO_CPU_VENDOR_INTEL) {
p2pLevel = PATH_PXB;
}
if (arch == NCCL_TOPO_CPU_ARCH_X86 && vendor == NCCL_TOPO_CPU_VENDOR_ZHAOXIN) {
p2pLevel = PATH_PXB;
}
compare:
// Compute the PCI distance and compare with the p2pLevel.
if (path->type <= p2pLevel) *p2p = 1;
#if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__)
#else
if (*p2p == 1) {
// NCCL_IGNORE_DISABLED_P2P=2 is used by unit tests that don't want to
// validate against NVML at all since they are pretending to be on other hw.
if (g1 != g2 && ncclParamIgnoreDisabledP2p() != 2) {
int indexes[3] = {-1,-1,-1};
int verticeN = 0;
NCCLCHECK(ncclNvmlEnsureInitialized());
indexes[verticeN++] = system->nodes[GPU].nodes[g1].gpu.dev;
if (intermediateIndex != -1) indexes[verticeN++] = system->nodes[GPU].nodes[intermediateIndex].gpu.dev;
indexes[verticeN++] = system->nodes[GPU].nodes[g2].gpu.dev;
for (int i=1; i < verticeN; i++) {
nvmlGpuP2PStatus_t status;
status = ncclNvmlDevicePairs[indexes[i-1]][indexes[i-0]].p2pStatusRead;
bool good = status == NVML_P2P_STATUS_OK;
status = ncclNvmlDevicePairs[indexes[i-1]][indexes[i-0]].p2pStatusWrite;
good &= status == NVML_P2P_STATUS_OK;
if (!good) {
if (!ncclParamIgnoreDisabledP2p()) {
if (path->type <= PATH_NVB) {
WARN("P2P is disabled between NVLINK connected GPUs %d and %d. This should not be the case given their connectivity, and is probably due to a hardware issue. If you still want to proceed, you can set NCCL_IGNORE_DISABLED_P2P=1.", indexes[i-1], indexes[i-0]);
return ncclUnhandledCudaError;
} else if (path->type < PATH_SYS) {
INFO(NCCL_INIT, "P2P is disabled between connected GPUs %d and %d. You can repress this message with NCCL_IGNORE_DISABLED_P2P=1.", indexes[i-1], indexes[i-0]);
}
}
*p2p = 0;
}
}
}
}
#endif
if (path->type == PATH_NVL) {
struct ncclTopoNode* gpu2 = system->nodes[GPU].nodes+g2;
// Enable P2P Read for Ampere/NVLink only
if (read && (gpu1->gpu.cudaCompCap == gpu2->gpu.cudaCompCap) && (gpu1->gpu.cudaCompCap == 80)) *read = 1;
}
return ncclSuccess;
}
// MNNVL: Check whether peers are in the same fabric cluster and clique
ncclResult_t ncclTopoCheckMNNVL(struct ncclTopoSystem* system, struct ncclPeerInfo* info1, struct ncclPeerInfo* info2, int* ret) {
*ret = 0;
nvmlGpuFabricInfoV_t *fabricInfo1 = &info1->fabricInfo;
nvmlGpuFabricInfoV_t *fabricInfo2 = &info2->fabricInfo;
// A zero UUID means we don't have MNNVL fabric info
if ((((long *)&fabricInfo2->clusterUuid)[0]|((long *)fabricInfo2->clusterUuid)[1]) == 0) return ncclSuccess;
if ((memcmp(fabricInfo1->clusterUuid, fabricInfo2->clusterUuid, NVML_GPU_FABRIC_UUID_LEN) == 0) &&
(fabricInfo1->cliqueId == fabricInfo2->cliqueId)) {
INFO(NCCL_NET, "MNNVL matching peer 0x%lx UUID %lx.%lx cliqueId 0x%x",
info2->busId, ((long *)fabricInfo2->clusterUuid)[0], ((long *)fabricInfo2->clusterUuid)[1], fabricInfo2->cliqueId);
*ret = 1;
}
return ncclSuccess;
}SHM：
看一下源码的逻辑过程和源码
1.初始化 ret 为0；
2.检查是否禁用共享内存SHM；
3.检查是否需要通过网络通信；
4.判断两个设备是否在同一主机；
5.检查设备是否共享相同的 /dev/shm；
6.满足条件，ret = 1；
/* Determine if two peers can communicate with SHM */
ncclResult_t shmCanConnect(int* ret, struct ncclTopoSystem* topo, struct ncclTopoGraph* graph, struct ncclPeerInfo* info1, struct ncclPeerInfo* info2) {
*ret = 0;
if (ncclParamShmDisable() == 1) return ncclSuccess;
int useNet = 0;
NCCLCHECK(ncclTopoCheckNet(topo, info1->busId, info2->busId, &useNet));
if (useNet) return ncclSuccess;
// Same host?
TRACE(NCCL_INIT|NCCL_SHM, "peer1 hostHash %lx peer2 hostHash %lx", info1->hostHash, info2->hostHash);
if (info1->hostHash != info2->hostHash) return ncclSuccess;
// Common /dev/shm (between containers) ?
TRACE(NCCL_INIT|NCCL_SHM, "peer1 shmDev %lx peer2 shmDev %lx", info1->shmDev, info2->shmDev);
if (info1->shmDev != info2->shmDev) return ncclSuccess;
*ret = 1;
return ncclSuccess;
}
