Add IPC teams

[ROCm/rocshmem commit: a59bdd4f6b]

projects/rocshmem/src/ipc/CMakeLists.txt

@@ -29,4 +29,5 @@ target_sources(
     context_ipc_device.cpp
     context_ipc_host.cpp
     backend_ipc.cpp
+    ipc_team.cpp
 )

projects/rocshmem/src/ipc/ipc_team.cpp

@@ -0,0 +1,56 @@
+/******************************************************************************
+ * Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in
+ * all copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
+ * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
+ * IN THE SOFTWARE.
+ *****************************************************************************/
+
+#include "ipc_team.hpp"
+
+#include "../backend_type.hpp"
+#include "backend_ipc.hpp"
+
+namespace rocshmem {
+
+IPCTeam::IPCTeam(Backend *backend, TeamInfo *team_info_parent,
+                     TeamInfo *team_info_world, int num_pes, int my_pe,
+                     MPI_Comm mpi_comm, int pool_index)
+    : Team(backend, team_info_parent, team_info_world, num_pes, my_pe,
+           mpi_comm) {
+  type = BackendType::IPC_BACKEND;
+  const IPCBackend *b = static_cast<const IPCBackend *>(backend);
+
+  pool_index_ = pool_index;
+
+  barrier_pSync =
+      &(b->barrier_pSync_pool[pool_index * ROC_SHMEM_BARRIER_SYNC_SIZE]);
+  reduce_pSync =
+      &(b->reduce_pSync_pool[pool_index * ROC_SHMEM_REDUCE_SYNC_SIZE]);
+  bcast_pSync = &(b->bcast_pSync_pool[pool_index * ROC_SHMEM_BCAST_SYNC_SIZE]);
+  alltoall_pSync =
+      &(b->alltoall_pSync_pool[pool_index * ROC_SHMEM_ALLTOALL_SYNC_SIZE]);
+
+  pWrk = reinterpret_cast<char *>(b->pWrk_pool) +
+         ROC_SHMEM_REDUCE_MIN_WRKDATA_SIZE * sizeof(double) * pool_index;
+  pAta = reinterpret_cast<char *>(b->pAta_pool) +
+         ROC_SHMEM_ATA_MAX_WRKDATA_SIZE * sizeof(double) * pool_index;
+}
+
+IPCTeam::~IPCTeam() {}
+
+}  // namespace rocshmem

projects/rocshmem/src/ipc/ipc_team.hpp

@@ -0,0 +1,50 @@
+/******************************************************************************
+ * Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in
+ * all copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
+ * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
+ * IN THE SOFTWARE.
+ *****************************************************************************/
+
+#ifndef LIBRARY_SRC_IPC_TEAM_HPP_
+#define LIBRARY_SRC_IPC_TEAM_HPP_
+
+#include "../team.hpp"
+
+namespace rocshmem {
+
+class IPCTeam : public Team {
+ public:
+  IPCTeam(Backend* handle, TeamInfo* team_info_wrt_parent,
+            TeamInfo* team_info_wrt_world, int num_pes, int my_pe,
+            MPI_Comm team_comm, int pool_index);
+
+  virtual ~IPCTeam();
+
+  long* barrier_pSync{nullptr};
+  long* reduce_pSync{nullptr};
+  long* bcast_pSync{nullptr};
+  long* alltoall_pSync{nullptr};
+  void* pWrk{nullptr};
+  void* pAta{nullptr};
+
+  int pool_index_{-1};
+};
+
+}  // namespace rocshmem
+
+#endif  // LIBRARY_SRC_IPC_TEAM_HPP_

projects/rocshmem/src/team.cpp

@@ -44,6 +44,10 @@ ROTeam* get_internal_ro_team(roc_shmem_team_t team) {
   return reinterpret_cast<ROTeam*>(team);
 }
 
+IPCTeam* get_internal_ipc_team(roc_shmem_team_t team) {
+  return reinterpret_cast<IPCTeam*>(team);
+}
+
 __host__ __device__ int team_translate_pe(roc_shmem_team_t src_team, int src_pe,
                                           roc_shmem_team_t dst_team) {
   if (src_team == ROC_SHMEM_TEAM_INVALID ||

projects/rocshmem/src/team.hpp

@@ -34,6 +34,7 @@ class Backend;
 class Team;
 class ROTeam;
 class GPUIBTeam;
+class IPCTeam;
 
 class TeamInfo {
  public:
@@ -162,6 +163,8 @@ GPUIBTeam* get_internal_gpu_ib_team(roc_shmem_team_t team);
 
 ROTeam* get_internal_ro_team(roc_shmem_team_t team);
 
+IPCTeam* get_internal_ipc_team(roc_shmem_team_t team);
+
 __host__ __device__ int team_translate_pe(roc_shmem_team_t src_team, int src_pe,
                                           roc_shmem_team_t dst_team);