diff --git a/CMakeLists.txt b/CMakeLists.txt index 433d0aefe3..c94f75514d 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -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) diff --git a/src/graph/rome_models.cc b/src/graph/rome_models.cc old mode 100755 new mode 100644 diff --git a/src/graph/rome_models.h b/src/graph/rome_models.h old mode 100755 new mode 100644 diff --git a/src/graph/xml.cc b/src/graph/xml.cc index add6b50409..124256f0d3 100644 --- a/src/graph/xml.cc +++ b/src/graph/xml.cc @@ -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 #include #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) diff --git a/src/misc/rocm_smi_wrap.cc b/src/misc/rocm_smi_wrap.cc new file mode 100644 index 0000000000..44b5fb153c --- /dev/null +++ b/src/misc/rocm_smi_wrap.cc @@ -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; +} diff --git a/src/transport/p2p.cc b/src/transport/p2p.cc index ea83e6975c..da594a7897 100644 --- a/src/transport/p2p.cc +++ b/src/transport/p2p.cc @@ -8,11 +8,8 @@ #include "comm.h" #include "graph.h" #include "utils.h" -#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__) -#include -#include -#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); }