AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Query XGMI link count through rocm_smi_lib API (#442)

Browse Source 
[ROCm/rccl commit: 14a184eb67]
This commit is contained in:
Wenkai Du
2021-10-26 10:30:20 -07:00
committed by GitHub
parent 0afd607328
commit 15143b1cfb
4 changed files with 67 additions and 56 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rccl/src/graph/xml.cc
+3 -3
View File
@@ -671,8 +671,8 @@ ncclResult_t ncclTopoGetXmlFromGpu(struct ncclXmlNode* pciNode, nvmlDevice_t nvm
for (int i=0; i<deviceCnt; i++) {
if (i != dev) {
RSMI_IO_LINK_TYPE rsmi_type;
int hops, bw;
if (rocm_smi_getLinkInfo(dev, i, &rsmi_type, &hops, &bw) == ncclSuccess) {
int hops, count;
if (rocm_smi_getLinkInfo(dev, i, &rsmi_type, &hops, &count) == ncclSuccess) {
if (rsmi_type == RSMI_IOLINK_TYPE_XGMI && hops == 1) {
char busIdStr[] = "00000000:00:00.0";
CUDACHECK(hipDeviceGetPCIBusId(busIdStr, sizeof(busIdStr), i));
@@ -685,7 +685,7 @@ ncclResult_t ncclTopoGetXmlFromGpu(struct ncclXmlNode* pciNode, nvmlDevice_t nvm
if (nvlNode == NULL) {
NCCLCHECK(xmlAddNode(xml, gpuNode, "xgmi", &nvlNode));
NCCLCHECK(xmlSetAttr(nvlNode, "target", lowerId));
NCCLCHECK(xmlSetAttrInt(nvlNode, "count", 1));
NCCLCHECK(xmlSetAttrInt(nvlNode, "count", count));
}
}
}
projects/rccl/src/include/rocm_smi_wrap.h
+2 -1
View File
@@ -24,10 +24,11 @@ THE SOFTWARE.
#define ROCM_SMI_WRAP_H_
#include "rocm_smi/rocm_smi.h"
#include "rocm_smi/rocm_smi64Config.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);
ncclResult_t rocm_smi_getLinkInfo(int srcDev, int dstDev, RSMI_IO_LINK_TYPE* rsmi_type, int *hops, int *count);
#endif
projects/rccl/src/misc/rocm_smi_wrap.cc
+14 -4
View File
@@ -74,7 +74,7 @@ ncclResult_t rocm_smi_getDeviceIndexByPciBusId(const char* pciBusId, uint32_t* d
}
}
ncclResult_t rocm_smi_getLinkInfo(int srcDev, int dstDev, RSMI_IO_LINK_TYPE* rsmi_type, int *hops, int *bw) {
ncclResult_t rocm_smi_getLinkInfo(int srcDev, int dstDev, RSMI_IO_LINK_TYPE* rsmi_type, int *hops, int *count) {
char srcStr[] = "00000000:00:00.0", dstStr[] = "00000000:00:00.0";
uint32_t srcIndex, dstIndex;
@@ -86,9 +86,19 @@ ncclResult_t rocm_smi_getLinkInfo(int srcDev, int dstDev, RSMI_IO_LINK_TYPE* rsm
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;
*count = 1;
if (*rsmi_type == RSMI_IOLINK_TYPE_XGMI && rsmi_weight == 15) {
*hops = 1;
#if rocm_smi_VERSION_MAJOR >= 5
uint64_t min_bw = 0, max_bw = 0;
rsmi_version_t version;
ROCMSMICHECK(rsmi_version_get(&version));
if (version.major >= 5)
ROCMSMICHECK(rsmi_minmax_bandwidth_get(srcIndex, dstIndex, &min_bw, &max_bw));
if (max_bw && min_bw)
*count = max_bw/min_bw;
#endif
}
return ncclSuccess;
}