Fix minor issues found during multi-node testing and addressing msccl failure

Author: Marzieh Berenjkoub

src/include/comm.h

@@ -737,6 +737,10 @@ struct ncclComm {
   bool useNetPXN;
   bool useGdr;
   int splitCount;
+  // gfx name from hipDeviceProp_t [RCCL]
+  char* archName;
+  // multiProcessorCount from hipDeviceProp_t [RCCL]
+  int cuCount;
 
   struct ncclDevrState devrState; // The symmetric runtime state
   struct ncclSymkState symkState; // The symmetric kernels state (built on previous)

src/include/msccl/msccl_kernel.h

@@ -9,7 +9,7 @@
 #define MSCCL_KERNEL_ENTRY_NAME(devredop, type, proto, fullOps) mscclKernel_##devredop##_##type##_##proto##_##fullOps
 
 #define MSCCL_DECL_KERNEL_ENTRY_FUNC_DEVREDOP_TYPE_PROTO(devredop, type, proto, fullOps) \
-__global__ void MSCCL_KERNEL_ENTRY_NAME(devredop, type, proto, fullOps)(struct ncclDevComm* comm, struct mscclAlgo* algo, struct mscclWork* work);
+__global__ void MSCCL_KERNEL_ENTRY_NAME(devredop, type, proto, fullOps)(struct ncclKernelComm* comm, struct mscclAlgo* algo, struct mscclWork* work);
 
 #define MSCCL_DECL_KERNEL_ENTRY_FUNC_DEVREDOP_TYPE(devredop, type, fullOps) \
   MSCCL_DECL_KERNEL_ENTRY_FUNC_DEVREDOP_TYPE_PROTO(devredop, type, LL, fullOps) \

src/misc/utils.cc

@@ -317,6 +317,16 @@ ncclResult_t ncclBitsToString(uint32_t bits, uint32_t mask, const char* (*toStr)
   return ncclSuccess;
 }
 
+size_t get_sc_page_size() {
+  static size_t cached_page_size = 0;
+  size_t ps = __atomic_load_n(&cached_page_size,__ATOMIC_RELAXED);
+  if (ps == 0) {
+      ps = (size_t)sysconf(_SC_PAGESIZE);
+      __atomic_store_n(&cached_page_size, ps,__ATOMIC_RELAXED);
+  }
+  return ps;
+}
+
 void get_aligned_ptr_and_size(const void *ptr, const size_t bufsize, void **aligned_ptr, size_t *aligned_size) {
   if (!aligned_ptr || !aligned_size) return;
   const size_t page_size = get_sc_page_size();