Use rocm_smi_lib for getting topology information (#402)
* Use rocm_smi_lib for getting topology information * Add rocm-smi-lib dependency to RCCL package
这个提交包含在：
Wenkai Du
GitHub
+5
-1
@@ -170,6 +170,7 @@ set(CC_SOURCES
|
||||
src/misc/utils.cc
|
||||
src/misc/ibvwrap.cc
|
||||
src/misc/nvmlwrap_stub.cc
|
||||
src/misc/rocm_smi_wrap.cc
|
||||
src/transport/coll_net.cc
|
||||
src/transport/net.cc
|
||||
src/transport/net_ib.cc
|
||||
@@ -240,7 +241,8 @@ if("${HIP_COMPILER}" MATCHES "hcc")
|
||||
endif()
|
||||
endif()
|
||||
|
||||
target_link_libraries(rccl PRIVATE hip::device dl)
|
||||
target_include_directories(rccl PRIVATE ${ROCM_PATH}/rocm_smi/include)
|
||||
target_link_libraries(rccl PRIVATE hip::device dl -lrocm_smi64 -L${ROCM_PATH}/rocm_smi/lib)
|
||||
target_link_libraries(rccl INTERFACE hip::host)
|
||||
|
||||
#Setup librccl.so version
|
||||
@@ -263,8 +265,10 @@ rocm_export_targets(NAMESPACE
|
||||
hip)
|
||||
|
||||
set(CPACK_DEBIAN_PACKAGE_DEPENDS "hip-rocclr (>= 3.5.0)")
|
||||
set(CPACK_DEBIAN_PACKAGE_DEPENDS "rocm-smi-lib (>= 4.0.0)")
|
||||
set(CPACK_DEBIAN_PACKAGE_SHLIBDEPS ON)
|
||||
set(CPACK_RPM_PACKAGE_REQUIRES "hip-rocclr >= 3.5.0")
|
||||
set(CPACK_RPM_PACKAGE_REQUIRES "rocm-smi-lib >= 4.0.0")
|
||||
set(CPACK_RPM_EXCLUDE_FROM_AUTO_FILELIST_ADDITION "/opt" "${ROCM_PATH}")
|
||||
|
||||
find_file (DEBIAN debian_version debconf.conf PATHS /etc)
|
||||
|
||||
可执行文件 -> 普通文件
可执行文件 -> 普通文件
+17
-3
@@ -17,6 +17,7 @@
|
||||
#include "core.h"
|
||||
#include "nvmlwrap.h"
|
||||
#include "xml.h"
|
||||
#include "rocm_smi_wrap.h"
|
||||
|
||||
/*******************/
|
||||
/* XML File Parser */
|
||||
@@ -656,9 +657,10 @@ ncclResult_t ncclTopoGetXmlFromGpu(struct ncclXmlNode* pciNode, nvmlDevice_t nvm
|
||||
CUDACHECK(hipGetDeviceCount(&deviceCnt));
|
||||
for (int i=0; i<deviceCnt; i++) {
|
||||
if (i != dev) {
|
||||
uint32_t link_type, hops;
|
||||
if (hipExtGetLinkTypeAndHopCount(dev, i, &link_type, &hops) == hipSuccess) {
|
||||
if (link_type == HSA_AMD_LINK_INFO_TYPE_XGMI && hops == 1) {
|
||||
RSMI_IO_LINK_TYPE rsmi_type;
|
||||
int hops, bw;
|
||||
if (rocm_smi_getLinkInfo(dev, i, &rsmi_type, &hops, &bw) == ncclSuccess) {
|
||||
if (rsmi_type == RSMI_IOLINK_TYPE_XGMI && hops == 1) {
|
||||
char busIdStr[] = "00000000:00:00.0";
|
||||
CUDACHECK(hipDeviceGetPCIBusId(busIdStr, sizeof(busIdStr), i));
|
||||
char lowerId[NVML_DEVICE_PCI_BUS_ID_BUFFER_SIZE];
|
||||
@@ -752,6 +754,17 @@ ncclResult_t ncclTopoFillGpu(struct ncclXml* xml, const char* busId, struct nccl
|
||||
struct ncclXmlNode* node;
|
||||
NCCLCHECK(ncclTopoGetPciNode(xml, busId, &node));
|
||||
NCCLCHECK(ncclTopoGetXmlFromSys(node, xml));
|
||||
#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__)
|
||||
uint32_t devIndex;
|
||||
static int rocmsmiInit = 0;
|
||||
if (rocmsmiInit == 0) {
|
||||
rocmsmiInit = (rocm_smi_init() != ncclSuccess) ? 2 : 1;
|
||||
}
|
||||
if (rocmsmiInit == 1) {
|
||||
if (rocm_smi_getDeviceIndexByPciBusId(busId, &devIndex) != ncclSuccess) devIndex = -1;
|
||||
}
|
||||
NCCLCHECK(ncclTopoGetXmlFromGpu(node, NULL, xml, gpuNode));
|
||||
#else
|
||||
nvmlDevice_t nvmlDev = NULL;
|
||||
static int nvmlInit = 0;
|
||||
if (nvmlInit == 0) {
|
||||
@@ -761,6 +774,7 @@ ncclResult_t ncclTopoFillGpu(struct ncclXml* xml, const char* busId, struct nccl
|
||||
if (wrapNvmlDeviceGetHandleByPciBusId(busId, &nvmlDev) != ncclSuccess) nvmlDev = NULL;
|
||||
}
|
||||
NCCLCHECK(ncclTopoGetXmlFromGpu(node, nvmlDev, xml, gpuNode));
|
||||
#endif
|
||||
return ncclSuccess;
|
||||
}
|
||||
|
||||
|
||||
@@ -0,0 +1,33 @@
|
||||
/*
|
||||
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#ifndef ROCM_SMI_WRAP_H_
|
||||
#define ROCM_SMI_WRAP_H_
|
||||
|
||||
#include "rocm_smi/rocm_smi.h"
|
||||
#include "nccl.h"
|
||||
|
||||
ncclResult_t rocm_smi_init();
|
||||
ncclResult_t rocm_smi_getDeviceIndexByPciBusId(const char* pciBusId, uint32_t* deviceIndex);
|
||||
ncclResult_t rocm_smi_getLinkInfo(int srcDev, int dstDev, RSMI_IO_LINK_TYPE* rsmi_type, int *hops, int *bw);
|
||||
|
||||
#endif
|
||||
+5
-3
@@ -28,6 +28,7 @@
|
||||
#include <sys/stat.h>
|
||||
#include <unistd.h>
|
||||
#include "graph/topo.h"
|
||||
#include "rocm_smi_wrap.h"
|
||||
|
||||
// [RCCL]
|
||||
#include "clique/CliqueManager.h"
|
||||
@@ -795,9 +796,10 @@ static ncclResult_t initTransportsRank(struct ncclComm* comm, ncclUniqueId* comm
|
||||
break;
|
||||
}
|
||||
|
||||
uint32_t linkType, hopCount;
|
||||
CUDACHECK(hipExtGetLinkTypeAndHopCount(i, j, &linkType, &hopCount));
|
||||
allXgmi &= (linkType == HSA_AMD_LINK_INFO_TYPE_XGMI);
|
||||
RSMI_IO_LINK_TYPE linkType;
|
||||
int hopCount, bw;
|
||||
NCCLCHECK(rocm_smi_getLinkInfo(i, j, &linkType, &hopCount, &bw));
|
||||
allXgmi &= (linkType == RSMI_IOLINK_TYPE_XGMI);
|
||||
}
|
||||
}
|
||||
if (hasPeerAccess)
|
||||
|
||||
@@ -0,0 +1,94 @@
|
||||
/*
|
||||
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
#include "rocm_smi_wrap.h"
|
||||
#include "core.h"
|
||||
#include "utils.h"
|
||||
|
||||
#define ROCMSMICHECK(cmd) do { \
|
||||
rsmi_status_t ret = cmd; \
|
||||
if( ret != RSMI_STATUS_SUCCESS ) { \
|
||||
const char *err; \
|
||||
rsmi_status_string(ret, &err); \
|
||||
WARN("ROCm SMI init failure %s", err); \
|
||||
return ncclInternalError; \
|
||||
} \
|
||||
} while(false)
|
||||
|
||||
|
||||
ncclResult_t rocm_smi_init() {
|
||||
ROCMSMICHECK(rsmi_init(0));
|
||||
rsmi_version_t version;
|
||||
ROCMSMICHECK(rsmi_version_get(&version));
|
||||
INFO(NCCL_INIT, "rocm_smi_lib: version %d.%d.%d.%s", version.major, version.minor, version.patch, version.build);
|
||||
return ncclSuccess;
|
||||
}
|
||||
|
||||
ncclResult_t rocm_smi_getDeviceIndexByPciBusId(const char* pciBusId, uint32_t* deviceIndex) {
|
||||
uint32_t i, num_devs = 0;
|
||||
int64_t busid;
|
||||
|
||||
busIdToInt64(pciBusId, &busid);
|
||||
/** convert to rocm_smi's bus ID format
|
||||
* | Name | Field |
|
||||
* ---------- | ------- |
|
||||
* | Domain | [64:32] |
|
||||
* | Reserved | [31:16] |
|
||||
* | Bus | [15: 8] |
|
||||
* | Device | [ 7: 3] |
|
||||
* | Function | [ 2: 0] |
|
||||
**/
|
||||
busid = ((busid&0xffff00000L)<<12)+((busid&0xff000L)>>4)+((busid&0xff0L)>>1)+(busid&0x7L);
|
||||
ROCMSMICHECK(rsmi_num_monitor_devices(&num_devs));
|
||||
for (i = 0; i < num_devs; i++) {
|
||||
uint64_t bdfid;
|
||||
ROCMSMICHECK(rsmi_dev_pci_id_get(i, &bdfid));
|
||||
if (bdfid == busid) break;
|
||||
}
|
||||
|
||||
if (i < num_devs) {
|
||||
*deviceIndex = i;
|
||||
return ncclSuccess;
|
||||
}
|
||||
else {
|
||||
WARN("rocm_smi_lib: %s device index not found", pciBusId);
|
||||
return ncclInternalError;
|
||||
}
|
||||
}
|
||||
|
||||
ncclResult_t rocm_smi_getLinkInfo(int srcDev, int dstDev, RSMI_IO_LINK_TYPE* rsmi_type, int *hops, int *bw) {
|
||||
char srcStr[] = "00000000:00:00.0", dstStr[] = "00000000:00:00.0";
|
||||
uint32_t srcIndex, dstIndex;
|
||||
|
||||
CUDACHECK(hipDeviceGetPCIBusId(srcStr, sizeof(srcStr), srcDev));
|
||||
CUDACHECK(hipDeviceGetPCIBusId(dstStr, sizeof(dstStr), dstDev));
|
||||
NCCLCHECK(rocm_smi_getDeviceIndexByPciBusId(srcStr, &srcIndex));
|
||||
NCCLCHECK(rocm_smi_getDeviceIndexByPciBusId(dstStr, &dstIndex));
|
||||
|
||||
uint64_t rsmi_hops, rsmi_weight;
|
||||
ROCMSMICHECK(rsmi_topo_get_link_type(srcIndex, dstIndex, &rsmi_hops, rsmi_type));
|
||||
ROCMSMICHECK(rsmi_topo_get_link_weight(srcIndex, dstIndex, &rsmi_weight));
|
||||
|
||||
*hops = 2;
|
||||
*bw = 0;
|
||||
if (*rsmi_type == RSMI_IOLINK_TYPE_XGMI && rsmi_weight == 15) *hops = 1;
|
||||
return ncclSuccess;
|
||||
}
|
||||
+5
-7
@@ -8,11 +8,8 @@
|
||||
#include "comm.h"
|
||||
#include "graph.h"
|
||||
#include "utils.h"
|
||||
#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__)
|
||||
#include <hsa/hsa.h>
|
||||
#include <hsa/hsa_ext_amd.h>
|
||||
#endif
|
||||
#include "bootstrap.h"
|
||||
#include "rocm_smi_wrap.h"
|
||||
|
||||
struct p2pConnectInfo {
|
||||
int rank;
|
||||
@@ -173,12 +170,13 @@ ncclResult_t p2pSendSetup(struct ncclComm* comm, struct ncclTopoGraph* graph, st
|
||||
NCCLCHECK(p2pGetInfo(comm->topo, myInfo, peerInfo, &useRead, &intermediateRank));
|
||||
|
||||
resources->next_hdp_reg = 0;
|
||||
uint32_t linktype, hops;
|
||||
if (hipExtGetLinkTypeAndHopCount(myInfo->cudaDev, peerInfo->cudaDev, &linktype, &hops) != hipSuccess) {
|
||||
RSMI_IO_LINK_TYPE linktype;
|
||||
int hops, bw;
|
||||
if (rocm_smi_getLinkInfo(myInfo->cudaDev, peerInfo->cudaDev, &linktype, &hops, &bw) != ncclSuccess) {
|
||||
INFO(NCCL_INIT|NCCL_P2P,"Ring %02d : %d -> %d failed to get link type and hop count", channelId, myInfo->rank, peerInfo->rank);
|
||||
return ncclInternalError;
|
||||
}
|
||||
if (linktype != HSA_AMD_LINK_INFO_TYPE_XGMI) {
|
||||
if (linktype != RSMI_IOLINK_TYPE_XGMI) {
|
||||
CUDACHECK(hipDeviceGetAttribute((int*)&resources->next_hdp_reg, hipDeviceAttributeHdpMemFlushCntl,peerInfo->cudaDev));
|
||||
TRACE(NCCL_INIT|NCCL_P2P,"Ring %02d : %d -> %d HDP %p", channelId, myInfo->rank, peerInfo->rank, resources->next_hdp_reg);
|
||||
}
|
||||
|
||||
在新工单中引用
屏蔽一个用户