Collect gcnArch and hipDeviceArch_t in XML (#252)

src/graph/topo.cc

@@ -307,6 +307,10 @@ ncclResult_t ncclTopoAddNic(struct ncclXmlNode* xmlNic, struct ncclTopoSystem* s
 
 ncclResult_t ncclTopoAddGpu(struct ncclXmlNode* xmlGpu, struct ncclTopoSystem* system, struct ncclTopoNode* gpu) {
   NCCLCHECK(xmlGetAttrInt(xmlGpu, "sm", &gpu->gpu.cudaCompCap));
+  NCCLCHECK(xmlGetAttrInt(xmlGpu, "gcn", &gpu->gpu.gcn));
+  rcclHipDeviceArch_t arch;
+  NCCLCHECK(xmlGetAttrInt(xmlGpu, "arch", &arch.value));
+  memcpy(&gpu->gpu.arch, &arch.arch, sizeof(hipDeviceArch_t));
   NCCLCHECK(xmlGetAttrInt(xmlGpu, "rank", &gpu->gpu.rank));
   NCCLCHECK(xmlGetAttrInt(xmlGpu, "dev", &gpu->gpu.dev));
   NCCLCHECK(xmlGetAttrInt(xmlGpu, "gdr", &gpu->gpu.gdrSupport));

src/graph/topo.h

@@ -90,6 +90,8 @@ struct ncclTopoNode {
       int rank;
       int cudaCompCap;
       int gdrSupport;
+      int gcn;
+      hipDeviceArch_t arch;
     }gpu;
     struct {
       uint64_t asic;

src/graph/xml.cc

@@ -599,6 +599,26 @@ ncclResult_t ncclTopoGetXmlFromGpu(struct ncclXmlNode* pciNode, nvmlDevice_t nvm
   int sm;
   NCCLCHECK(xmlGetAttrInt(gpuNode, "sm", &sm));
 
+  int gcn;
+  NCCLCHECK(xmlGetAttrIndex(gpuNode, "gcn", &index));
+  if (index == -1) {
+    hipDeviceProp_t devProp;
+    CUDACHECK(hipGetDeviceProperties(&devProp, dev));
+    gcn = devProp.gcnArch;
+    NCCLCHECK(xmlSetAttrInt(gpuNode, "gcn", gcn));
+  }
+  NCCLCHECK(xmlGetAttrInt(gpuNode, "gcn", &gcn));
+
+  rcclHipDeviceArch_t arch;
+  NCCLCHECK(xmlGetAttrIndex(gpuNode, "arch", &index));
+  if (index == -1) {
+    hipDeviceProp_t devProp;
+    CUDACHECK(hipGetDeviceProperties(&devProp, dev));
+    memcpy(&arch.arch, &devProp.arch, sizeof(hipDeviceArch_t));
+    NCCLCHECK(xmlSetAttrInt(gpuNode, "arch", arch.value));
+  }
+  NCCLCHECK(xmlGetAttrInt(gpuNode, "arch", &arch.value));
+
   struct ncclXmlNode* nvlNode = NULL;
   NCCLCHECK(xmlGetSub(pciNode, "nvlink", &nvlNode));
   if (nvlNode == NULL) {

src/graph/xml.h

@@ -37,7 +37,7 @@ struct ncclXml {
 };
 
 /* File functions */
-#define NCCL_TOPO_XML_VERSION 1
+#define NCCL_TOPO_XML_VERSION 2
 ncclResult_t ncclTopoGetXmlFromFile(const char* xmlTopoFile, struct ncclXml* xml);
 ncclResult_t ncclTopoDumpXmlToFile(const char* xmlTopoFile, struct ncclXml* xml);
 #define NCCL_GRAPH_XML_VERSION 1
@@ -235,4 +235,10 @@ static ncclResult_t kvConvertToStr(int value, const char** str, struct kvDict* d
   return ncclInternalError;
 }
 
+typedef union {
+  hipDeviceArch_t arch;
+  int value;
+  static_assert(sizeof(hipDeviceArch_t) == sizeof(int),
+      "value must be the same size of hipDeviceArch_t.");
+} rcclHipDeviceArch_t;
 #endif