From 6f8d2df73a3af37f62381e0438e146587d93885d Mon Sep 17 00:00:00 2001 From: avinashkethineedi Date: Thu, 22 Aug 2024 09:15:44 -0700 Subject: [PATCH] Add IPC teams [ROCm/rocshmem commit: a59bdd4f6b6792a097342e774e929fa26000f483] --- projects/rocshmem/src/ipc/CMakeLists.txt | 1 + projects/rocshmem/src/ipc/ipc_team.cpp | 56 ++++++++++++++++++++++++ projects/rocshmem/src/ipc/ipc_team.hpp | 50 +++++++++++++++++++++ projects/rocshmem/src/team.cpp | 4 ++ projects/rocshmem/src/team.hpp | 3 ++ 5 files changed, 114 insertions(+) create mode 100644 projects/rocshmem/src/ipc/ipc_team.cpp create mode 100644 projects/rocshmem/src/ipc/ipc_team.hpp diff --git a/projects/rocshmem/src/ipc/CMakeLists.txt b/projects/rocshmem/src/ipc/CMakeLists.txt index b8d50d12e2..e8d64ff523 100644 --- a/projects/rocshmem/src/ipc/CMakeLists.txt +++ b/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 ) diff --git a/projects/rocshmem/src/ipc/ipc_team.cpp b/projects/rocshmem/src/ipc/ipc_team.cpp new file mode 100644 index 0000000000..433ca32d4f --- /dev/null +++ b/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(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(b->pWrk_pool) + + ROC_SHMEM_REDUCE_MIN_WRKDATA_SIZE * sizeof(double) * pool_index; + pAta = reinterpret_cast(b->pAta_pool) + + ROC_SHMEM_ATA_MAX_WRKDATA_SIZE * sizeof(double) * pool_index; +} + +IPCTeam::~IPCTeam() {} + +} // namespace rocshmem diff --git a/projects/rocshmem/src/ipc/ipc_team.hpp b/projects/rocshmem/src/ipc/ipc_team.hpp new file mode 100644 index 0000000000..e931adcfe4 --- /dev/null +++ b/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_ diff --git a/projects/rocshmem/src/team.cpp b/projects/rocshmem/src/team.cpp index fefb520328..21ed00037c 100644 --- a/projects/rocshmem/src/team.cpp +++ b/projects/rocshmem/src/team.cpp @@ -44,6 +44,10 @@ ROTeam* get_internal_ro_team(roc_shmem_team_t team) { return reinterpret_cast(team); } +IPCTeam* get_internal_ipc_team(roc_shmem_team_t team) { + return reinterpret_cast(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 || diff --git a/projects/rocshmem/src/team.hpp b/projects/rocshmem/src/team.hpp index 56475e799b..0dfaac6ae3 100644 --- a/projects/rocshmem/src/team.hpp +++ b/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);