Enable GDA+IPC (#249)

* Enable GDA+IPC
Fix ROCSHMEM_DISABLE_IPC for both RO and GDA

* add more functionality to bootstrap class

we need a few more functions in the boostrap class to be able to fully
handle the rocshmem requirements:
 - add a function to return the list of local ranks
 - provide a groupAllgather operation which takes a vector of ranks
   participating
 - provide a groupAlltoall operation which takes a vector of ranks
   participating

Also, update the functionality of the gda-Alltoall and gda-Allreduce
operations to take advantage of these functions.

* ipc_policy adapted to use bootstrap groupallgather

* bugfix: there was a mistake in computing sendto in groupallgather

* bugfix: shm_size and shm_rank were set in a local variable rather than
the class member

* mpi-bootstrap: remove an unecessary allgather

---------

Co-authored-by: Edgar Gabriel <Edgar.Gabriel@amd.com>

src/gda/context_gda_host.cpp

@@ -42,9 +42,29 @@ __host__ GDAHostContext::GDAHostContext(Backend *backend,
   host_interface = b->host_interface;
 
   context_window_info = host_interface->acquire_window_context();
+
+  int *pes_with_ipc_avail = new int[backend->ipcImpl.shm_size];
+  char** ipc_bases = new char*[b->ipcImpl.shm_size];
+  if (backend->ipcImpl.pes_with_ipc_avail != nullptr) {
+    CHECK_HIP(hipMemcpy(pes_with_ipc_avail,
+                  backend->ipcImpl.pes_with_ipc_avail,
+                  backend->ipcImpl.shm_size * sizeof(int),
+                  hipMemcpyDeviceToHost));
+    CHECK_HIP(hipMemcpy(ipc_bases,
+                  backend->ipcImpl.ipc_bases,
+                  backend->ipcImpl.shm_size * sizeof(char *),
+                  hipMemcpyDeviceToHost));
+  }
+  ipcImpl_.pes_with_ipc_avail = pes_with_ipc_avail;
+  ipcImpl_.ipc_bases = ipc_bases;
+  ipcImpl_.shm_size = backend->ipcImpl.shm_size;
+  ipcImpl_.shm_rank = backend->ipcImpl.shm_rank;
 }
 
 __host__ GDAHostContext::~GDAHostContext() {
+  delete[] ipcImpl_.pes_with_ipc_avail;
+  delete[] ipcImpl_.ipc_bases;
+
   host_interface->release_window_context(context_window_info);
 }
 
@@ -78,8 +98,12 @@ __host__ void GDAHostContext::quiet() {
 
 __host__ void *GDAHostContext::shmem_ptr(const void *dest, int pe) {
   void *ret = nullptr;
-  //not implemented, returning nullptr is spec-valid
-  //TODO: copy ipc handover from RO when IPC+GDA is implemented
+  int local_pe{-1};
+  if (ipcImpl_.isIpcAvailable(my_pe, pe, &local_pe)) {
+    void *dst = const_cast<void *>(dest);
+    uint64_t L_offset = reinterpret_cast<char *>(dst) - ipcImpl_.ipc_bases[ipcImpl_.shm_rank];
+    ret = ipcImpl_.ipc_bases[local_pe] + L_offset;
+  }
   return ret;
 }