From 84f8ba3bb06895169b144c21c4c5bc47f3c73709 Mon Sep 17 00:00:00 2001
From: Wenkai Du <43822138+wenkaidu@users.noreply.github.com>
Date: Wed, 24 Jun 2020 11:25:55 -0700
Subject: [PATCH] Revert use posix_memalign for network buffer allocation on
 host memory (#222)

---
 CMakeLists.txt        |  2 +-
 src/include/devcomm.h |  3 ---
 src/transport/net.cc  | 36 ------------------------------------
 3 files changed, 1 insertion(+), 40 deletions(-)

diff --git a/CMakeLists.txt b/CMakeLists.txt
index ff79e764ba..fa42e4c6cd 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -169,7 +169,7 @@ endforeach()
 
 if("${HIP_COMPILER}" MATCHES "clang")
   foreach(target ${AMDGPU_TARGETS})
-    target_compile_options(rccl PRIVATE --amdgpu-target=${target} PRIVATE -fgpu-rdc -DUSE_MEMALIGN)
+    target_compile_options(rccl PRIVATE --amdgpu-target=${target} PRIVATE -fgpu-rdc)
   endforeach()
   target_link_libraries(rccl PRIVATE -fgpu-rdc)
   target_include_directories(rccl PRIVATE /opt/rocm/hsa/include)
diff --git a/src/include/devcomm.h b/src/include/devcomm.h
index e136444f92..e5c5f43727 100644
--- a/src/include/devcomm.h
+++ b/src/include/devcomm.h
@@ -92,9 +92,6 @@ static_assert(NCCL_LL_CLEAN_MASK % NCCL_STEPS == 0, "Invalid NCCL_LL_CLEAN_MASK
 struct ncclConnInfo {
   // Regular comm mechanism
   char *buffs[NCCL_NUM_PROTOCOLS]; // Local for recv, remote for send
-#ifdef USE_MEMALIGN
-  char *hostBuffs[NCCL_NUM_PROTOCOLS]; // Local for recv, remote for send
-#endif
   uint64_t *tail;     // Local for recv, remote for send
   uint64_t *head;     // Local for send, remote for recv
   uint64_t *opCountLoc; // opCount of local rank
diff --git a/src/transport/net.cc b/src/transport/net.cc
index 73e9bfffc1..7a9bf960a5 100644
--- a/src/transport/net.cc
+++ b/src/transport/net.cc
@@ -26,9 +26,6 @@ struct netSendResources {
   int netDev;
   int useGdr;
   char* buffers[LOC_COUNT];
-#ifdef USE_MEMALIGN
-  char* hostBuffer;
-#endif
   int buffSizes[LOC_COUNT];
   void* mhandles[LOC_COUNT];
   void** mhandlesProto[NCCL_NUM_PROTOCOLS];
@@ -96,24 +93,10 @@ ncclResult_t netSendSetup(struct ncclTopoSystem* topo, struct ncclTopoGraph* gra
   }
   char line[16];
   if (resources->buffSizes[LOC_HOSTMEM]) {
-#ifdef USE_MEMALIGN
-    int page_size = getpagesize();
-    if (posix_memalign((void **)&resources->hostBuffer, page_size, resources->buffSizes[LOC_HOSTMEM])) {
-      WARN("Failed to posix_memalign %d bytes", resources->buffSizes[LOC_HOSTMEM]);
-      return ncclSystemError;
-    }
-    CUDACHECK(hipHostRegister(resources->hostBuffer, resources->buffSizes[LOC_HOSTMEM], hipHostRegisterMapped));
-    CUDACHECK(hipHostGetDevicePointer((void **)resources->buffers+LOC_HOSTMEM, resources->hostBuffer, 0));
-#else
     NCCLCHECK(ncclCudaHostCalloc(resources->buffers+LOC_HOSTMEM, resources->buffSizes[LOC_HOSTMEM]));
-#endif
     int status[1] = {-1};
     line[0]= 0;
-#ifdef USE_MEMALIGN
-    if (!move_pages(0, 1, (void **)&resources->hostBuffer, NULL, status, 0))
-#else
     if (!move_pages(0, 1, (void **)resources->buffers+LOC_HOSTMEM, NULL, status, 0))
-#endif
       sprintf(line, "/MEM%d", status[0]);
   }
 
@@ -122,12 +105,6 @@ ncclResult_t netSendSetup(struct ncclTopoSystem* topo, struct ncclTopoGraph* gra
   for (int p=0; p<NCCL_NUM_PROTOCOLS; p++) {
     resources->mhandlesProto[p] = resources->mhandles+protoLoc[p];
     send->conn.buffs[p] = resources->buffers[protoLoc[p]] + offsets[protoLoc[p]];
-#ifdef USE_MEMALIGN
-    if (protoLoc[p] == LOC_HOSTMEM)
-      send->conn.hostBuffs[p] = resources->hostBuffer + offsets[protoLoc[p]];
-    else
-      send->conn.hostBuffs[p] = send->conn.buffs[p];
-#endif
     offsets[protoLoc[p]] += buffSizes[p];
   }
 
@@ -205,11 +182,7 @@ ncclResult_t netSendConnect(struct ncclConnect* connectInfo, int nranks, int ran
     NCCLCHECK(ncclNetRegMr(resources->netSendComm, resources->buffers[LOC_DEVMEM], resources->buffSizes[LOC_DEVMEM], NCCL_PTR_CUDA, &resources->mhandles[LOC_DEVMEM]));
   }
   if (resources->buffSizes[LOC_HOSTMEM]) {
-#ifdef USE_MEMALIGN
-    NCCLCHECK(ncclNetRegMr(resources->netSendComm, resources->hostBuffer, resources->buffSizes[LOC_HOSTMEM], NCCL_PTR_HOST, &resources->mhandles[LOC_HOSTMEM]));
-#else
     NCCLCHECK(ncclNetRegMr(resources->netSendComm, resources->buffers[LOC_HOSTMEM], resources->buffSizes[LOC_HOSTMEM], NCCL_PTR_HOST, &resources->mhandles[LOC_HOSTMEM]));
-#endif
   }
   return ncclSuccess;
 }
@@ -240,12 +213,7 @@ ncclResult_t netSendFree(void* transportResources) {
     if (resources->buffers[l])
       NCCLCHECK(ncclNetDeregMr(resources->netSendComm, resources->mhandles[l]));
   }
-#ifdef USE_MEMALIGN
-  CUDACHECK(hipHostUnregister(resources->hostBuffer));
-  free(resources->hostBuffer);
-#else
   NCCLCHECK(ncclCudaHostFree(resources->buffers[LOC_HOSTMEM]));
-#endif
   CUDACHECK(hipFree(resources->buffers[LOC_DEVMEM]));
   NCCLCHECK(ncclNetCloseSend(resources->netSendComm));
   free(resources);
@@ -283,11 +251,7 @@ ncclResult_t netSendProxy(struct ncclProxyArgs* args) {
   if (args->state == ncclProxyOpProgress) {
     int p = args->protocol;
     int stepSize = args->connector->comm->buffSizes[p] / NCCL_STEPS;
-#ifdef USE_MEMALIGN
-    char* localBuff = args->connector->conn.hostBuffs[p];
-#else
     char* localBuff = args->connector->conn.buffs[p];
-#endif
     void* mhandle = *(resources->mhandlesProto[p]);
     args->idle = 1;
     if (args->head < args->end) {