diff --git a/CMakeLists.txt b/CMakeLists.txt index 8603737308..f9c3d83d10 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -54,7 +54,8 @@ string(REGEX REPLACE "([0-9]+)\.([0-9]+)\.([0-9]+)(.*)" "\\1.\\2.\\3" ROCSHMEM_V ############################################################################### option(DEBUG "Enable debug trace" OFF) option(PROFILE "Enable statistics and timing support" OFF) -option(USE_GPU_IB "Enable GPU_IB conduit. If off, RO_NET will be used" ON) +option(USE_GPU_IB "Enable GPU_IB conduit." ON) +option(USE_RO "Enable RO conduit." ON) option(USE_DC "Enable IB dynamically connected transport (DC)" OFF) option(USE_IPC "Enable IPC support (using HIP)" OFF) option(USE_THREADS "Enable workgroup threads to share network queues" OFF) diff --git a/cmake/config.h.in b/cmake/config.h.in index f409c10ba3..8067651a01 100644 --- a/cmake/config.h.in +++ b/cmake/config.h.in @@ -1,6 +1,7 @@ #cmakedefine DEBUG #cmakedefine PROFILE #cmakedefine USE_GPU_IB +#cmakedefine USE_RO #cmakedefine USE_DC #cmakedefine USE_IPC #cmakedefine USE_THREADS diff --git a/scripts/build_configs/ipc_single b/scripts/build_configs/ipc_single new file mode 100755 index 0000000000..2013b2b67c --- /dev/null +++ b/scripts/build_configs/ipc_single @@ -0,0 +1,30 @@ +#!/bin/bash +# Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. + +if [ -z $1 ] +then + install_path=~/rocshmem +else + install_path=$1 +fi + +src_path=$(dirname "$(realpath $0)")/../../ + +cmake \ + -DCMAKE_BUILD_TYPE=Release \ + -DCMAKE_INSTALL_PREFIX=$install_path \ + -DCMAKE_VERBOSE_MAKEFILE=OFF \ + -DDEBUG=OFF \ + -DPROFILE=OFF \ + -DUSE_GPU_IB=OFF \ + -DUSE_RO=OFF \ + -DUSE_DC=OFF \ + -DUSE_IPC=ON \ + -DUSE_COHERENT_HEAP=ON \ + -DUSE_THREADS=OFF \ + -DUSE_WF_COAL=OFF \ + -DUSE_SINGLE_NODE=ON \ + -DUSE_HOST_SIDE_HDP_FLUSH=OFF\ + $src_path +cmake --build . --parallel 8 +cmake --install . \ No newline at end of file diff --git a/scripts/build_configs/ro_net b/scripts/build_configs/ro_net index 30c4ad27bc..17809fa0a9 100755 --- a/scripts/build_configs/ro_net +++ b/scripts/build_configs/ro_net @@ -18,7 +18,7 @@ cmake \ -DPROFILE=OFF \ -DUSE_GPU_IB=OFF \ -DUSE_DC=OFF \ - -DUSE_IPC=OFF \ + -DUSE_IPC=ON \ -DUSE_THREADS=ON \ -DUSE_WF_COAL=OFF \ -DUSE_COHERENT_HEAP=ON \ diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 823102d8a1..5c31872765 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -30,7 +30,6 @@ target_sources( backend_bc.cpp context_host.cpp context_device.cpp - ipc_policy.cpp mpi_init_singleton.cpp roc_shmem_gpu.cpp roc_shmem.cpp @@ -38,6 +37,7 @@ target_sources( team_tracker.cpp util.cpp wf_coal_policy.cpp + ipc_policy.cpp ) target_compile_options( @@ -60,8 +60,10 @@ target_compile_options( ############################################################################### IF (USE_GPU_IB) add_subdirectory(gpu_ib) -ELSE() +ELSEIF(USE_RO) add_subdirectory(reverse_offload) +ELSE() +add_subdirectory(ipc) ENDIF() add_subdirectory(containers) add_subdirectory(host) diff --git a/src/backend_bc.cpp b/src/backend_bc.cpp index 1c3a619887..c55b1cd2d1 100644 --- a/src/backend_bc.cpp +++ b/src/backend_bc.cpp @@ -25,10 +25,12 @@ #include "backend_type.hpp" #include "context_incl.hpp" -#ifndef USE_GPU_IB +#ifdef USE_GPU_IB +#include "gpu_ib/backend_ib.hpp" +#elif defined(USE_RO) #include "reverse_offload/backend_ro.hpp" #else -#include "gpu_ib/backend_ib.hpp" +#include "ipc/backend_ipc.hpp" #endif namespace rocshmem { @@ -201,18 +203,22 @@ void Backend::reset_stats() { } __device__ bool Backend::create_ctx(int64_t option, roc_shmem_ctx_t* ctx) { -#ifndef USE_GPU_IB +#ifdef USE_GPU_IB + return static_cast(this)->create_ctx(option, ctx); +#elif defined(USE_RO) return static_cast(this)->create_ctx(option, ctx); #else - return static_cast(this)->create_ctx(option, ctx); + return static_cast(this)->create_ctx(option, ctx); #endif } __device__ void Backend::destroy_ctx(roc_shmem_ctx_t* ctx) { -#ifndef USE_GPU_IB +#ifdef USE_GPU_IB + static_cast(this)->destroy_ctx(ctx); +#elif defined(USE_RO) static_cast(this)->destroy_ctx(ctx); #else - static_cast(this)->destroy_ctx(ctx); + static_cast(this)->destroy_ctx(ctx); #endif } diff --git a/src/backend_type.hpp b/src/backend_type.hpp index ce6ccb64cc..50818eb178 100644 --- a/src/backend_type.hpp +++ b/src/backend_type.hpp @@ -44,7 +44,7 @@ namespace rocshmem { * @note Derived classes which use Backend as a base class must add * themselves to this enum class to support static polymorphism. */ -enum class BackendType { RO_BACKEND, GPU_IB_BACKEND }; +enum class BackendType { RO_BACKEND, GPU_IB_BACKEND, IPC_BACKEND }; /** * @brief Helper macro for some dispatch calls @@ -57,9 +57,12 @@ enum class BackendType { RO_BACKEND, GPU_IB_BACKEND }; #ifdef USE_GPU_IB #define DISPATCH(Func) \ static_cast(this)->Func; -#else +#elif defined(USE_RO) #define DISPATCH(Func) \ static_cast(this)->Func; +#else +#define DISPATCH(Func) \ + static_cast(this)->Func; #endif /** @@ -69,10 +72,15 @@ enum class BackendType { RO_BACKEND, GPU_IB_BACKEND }; #define DISPATCH_RET(Func) \ auto ret_val = static_cast(this)->Func; \ return ret_val; -#else +#elif defined(USE_RO) #define DISPATCH_RET(Func) \ auto ret_val = static_cast(this)->Func; \ return ret_val; +#else +#define DISPATCH_RET(Func) \ + auto ret_val{0}; \ + ret_val = static_cast(this)->Func; \ + return ret_val; #endif /** * @brief Device static dispatch method call with a return type of pointer. @@ -82,11 +90,16 @@ enum class BackendType { RO_BACKEND, GPU_IB_BACKEND }; void *ret_val{nullptr}; \ ret_val = static_cast(this)->Func; \ return ret_val; -#else +#elif defined(USE_RO) #define DISPATCH_RET_PTR(Func) \ void *ret_val{nullptr}; \ ret_val = static_cast(this)->Func; \ return ret_val; +#else +#define DISPATCH_RET_PTR(Func) \ + void *ret_val{nullptr}; \ + ret_val = static_cast(this)->Func; \ + return ret_val; #endif /** @@ -98,8 +111,10 @@ enum class BackendType { RO_BACKEND, GPU_IB_BACKEND }; */ #ifdef USE_GPU_IB #define HOST_DISPATCH(Func) static_cast(this)->Func; -#else +#elif defined(USE_RO) #define HOST_DISPATCH(Func) static_cast(this)->Func; +#else +#define HOST_DISPATCH(Func) static_cast(this)->Func; #endif /** * @brief Host static dispatch method call with return value. @@ -113,10 +128,15 @@ enum class BackendType { RO_BACKEND, GPU_IB_BACKEND }; #define HOST_DISPATCH_RET(Func) \ auto ret_val = static_cast(this)->Func; \ return ret_val; -#else +#elif defined(USE_RO) #define HOST_DISPATCH_RET(Func) \ auto ret_val = static_cast(this)->Func; \ return ret_val; +#else +#define HOST_DISPATCH_RET(Func) \ + auto ret_val{0}; \ + ret_val = static_cast(this)->Func; \ + return ret_val; #endif } // namespace rocshmem diff --git a/src/context_incl.hpp b/src/context_incl.hpp index ea39a2ead5..f86ca3dfe8 100644 --- a/src/context_incl.hpp +++ b/src/context_incl.hpp @@ -29,9 +29,12 @@ #ifdef USE_GPU_IB #include "gpu_ib/context_ib_device.hpp" #include "gpu_ib/context_ib_host.hpp" -#else +#elif defined (USE_RO) #include "reverse_offload/context_ro_device.hpp" #include "reverse_offload/context_ro_host.hpp" +#else +#include "ipc/context_ipc_device.hpp" +#include "ipc/context_ipc_host.hpp" #endif #endif // LIBRARY_SRC_CONTEXT_INCL_HPP_ diff --git a/src/context_tmpl_device.hpp b/src/context_tmpl_device.hpp index 213fb90d3b..d6d0df63fa 100644 --- a/src/context_tmpl_device.hpp +++ b/src/context_tmpl_device.hpp @@ -27,8 +27,10 @@ #include "backend_type.hpp" #ifdef USE_GPU_IB #include "gpu_ib/context_ib_device.hpp" -#else +#elif defined(USE_RO) #include "reverse_offload/context_ro_device.hpp" +#else +#include "ipc/context_ipc_device.hpp" #endif namespace rocshmem { diff --git a/src/context_tmpl_host.hpp b/src/context_tmpl_host.hpp index 56b649e19d..bce99c1d6c 100644 --- a/src/context_tmpl_host.hpp +++ b/src/context_tmpl_host.hpp @@ -27,8 +27,10 @@ #include "backend_type.hpp" #ifdef USE_GPU_IB #include "gpu_ib/context_ib_host.hpp" -#else +#elif defined(USE_RO) #include "reverse_offload/context_ro_host.hpp" +#else +#include "ipc/context_ipc_host.hpp" #endif namespace rocshmem { diff --git a/src/ipc/CMakeLists.txt b/src/ipc/CMakeLists.txt new file mode 100644 index 0000000000..07a40538a5 --- /dev/null +++ b/src/ipc/CMakeLists.txt @@ -0,0 +1,34 @@ +############################################################################### +# 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. +############################################################################### + +############################################################################### +# ADD ROCSHMEM TARGET FOR FILES IN CURRENT DIRECTORY +############################################################################### +target_sources( + ${PROJECT_NAME} + PRIVATE + context_ipc_device.cpp + context_ipc_host.cpp + backend_ipc.cpp + ipc_team.cpp + context_ipc_device_coll.cpp +) diff --git a/src/ipc/backend_ipc.cpp b/src/ipc/backend_ipc.cpp new file mode 100644 index 0000000000..58caaf53f7 --- /dev/null +++ b/src/ipc/backend_ipc.cpp @@ -0,0 +1,392 @@ +/****************************************************************************** + * 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 "backend_ipc.hpp" +#include "ipc_team.hpp" + +namespace rocshmem { + +#define NET_CHECK(cmd) \ + { \ + if (cmd != MPI_SUCCESS) { \ + fprintf(stderr, "Unrecoverable error: MPI Failure\n"); \ + abort() ; \ + } \ + } + +extern roc_shmem_ctx_t ROC_SHMEM_HOST_CTX_DEFAULT; + +roc_shmem_team_t get_external_team(GPUIBTeam *team) { + return reinterpret_cast(team); +} + +int get_ls_non_zero_bit(char *bitmask, int mask_length) { + int position = -1; + + for (int bit_i = 0; bit_i < mask_length; bit_i++) { + int byte_i = bit_i / CHAR_BIT; + if (bitmask[byte_i] & (1 << (bit_i % CHAR_BIT))) { + position = bit_i; + break; + } + } + + return position; +} + +IPCBackend::IPCBackend(MPI_Comm comm) + : Backend() { + type = BackendType::IPC_BACKEND; + + if (auto maximum_num_contexts_str = getenv("ROC_SHMEM_MAX_NUM_CONTEXTS")) { + std::stringstream sstream(maximum_num_contexts_str); + sstream >> maximum_num_contexts_; + } + + init_mpi_once(comm); + + initIPC(); + + auto *bp{ipc_backend_proxy.get()}; + + bp->heap_ptr = &heap; + + /* Initialize the host interface */ + host_interface = + new HostInterface(hdp_proxy_.get(), thread_comm, &heap); + + default_host_ctx = std::make_unique(this, 0); + + ROC_SHMEM_HOST_CTX_DEFAULT.ctx_opaque = default_host_ctx.get(); + + init_g_ret(&heap, thread_comm, MAX_NUM_BLOCKS, &bp->g_ret); + + allocate_atomic_region(&bp->atomic_ret, MAX_NUM_BLOCKS); + + default_context_proxy_ = IPCDefaultContextProxyT(this); + + setup_team_world(); + + roc_shmem_collective_init(); + + teams_init(); + + setup_ctxs(); + +} + +IPCBackend::~IPCBackend() { + /* + * Validate that a handle was passed that is not a nullptr. + */ + auto *bp{ipc_backend_proxy.get()}; + assert(bp); + + /* + * Free the atomic_ret array. + */ + CHECK_HIP(hipFree(bp->atomic_ret->atomic_base_ptr)); + + // TODO(Avinash) Free g_ret + + // delete host_interface; + // host_interface = nullptr; + + /** + * Destroy teams infrastructure + * and team world + */ + teams_destroy(); + auto *team_world{team_tracker.get_team_world()}; + team_world->~Team(); + CHECK_HIP(hipFree(team_world)); + + CHECK_HIP(hipFree(ctx_array)); +} + +void IPCBackend::setup_ctxs() { + CHECK_HIP(hipMalloc(&ctx_array, sizeof(IPCContext) * maximum_num_contexts_)); + for (int i = 0; i < maximum_num_contexts_; i++) { + new (&ctx_array[i]) IPCContext(this); + ctx_free_list.get()->push_back(ctx_array + i); + } +} + +__device__ bool IPCBackend::create_ctx(int64_t options, roc_shmem_ctx_t *ctx) { + IPCContext *ctx_{nullptr}; + + auto pop_result = ctx_free_list.get()->pop_front(); + if (!pop_result.success) { + return false; + } + ctx_ = pop_result.value; + + ctx->ctx_opaque = ctx_; + return true; +} + +__device__ void IPCBackend::destroy_ctx(roc_shmem_ctx_t *ctx) { + ctx_free_list.get()->push_back(static_cast(ctx->ctx_opaque)); +} + +void IPCBackend::setup_team_world() { + TeamInfo *team_info_wrt_parent, *team_info_wrt_world; + + /** + * Allocate device-side memory for team_world and construct a + * IPC team in it. + */ + CHECK_HIP(hipMalloc(&team_info_wrt_parent, sizeof(TeamInfo))); + CHECK_HIP(hipMalloc(&team_info_wrt_world, sizeof(TeamInfo))); + + new (team_info_wrt_parent) TeamInfo(nullptr, 0, 1, num_pes); + new (team_info_wrt_world) TeamInfo(nullptr, 0, 1, num_pes); + + IPCTeam *team_world{nullptr}; + CHECK_HIP(hipMalloc(&team_world, sizeof(IPCTeam))); + new (team_world) IPCTeam(this, team_info_wrt_parent, team_info_wrt_world, + num_pes, my_pe, thread_comm, 0); + team_tracker.set_team_world(team_world); + + /** + * Copy the address to ROC_SHMEM_TEAM_WORLD. + */ + ROC_SHMEM_TEAM_WORLD = reinterpret_cast(team_world); +} + +void IPCBackend::init_mpi_once(MPI_Comm comm) { + int init_done{}; + NET_CHECK(MPI_Initialized(&init_done)); + + int provided{}; + if (!init_done) { + NET_CHECK(MPI_Init_thread(0, 0, MPI_THREAD_MULTIPLE, &provided)); + if (provided != MPI_THREAD_MULTIPLE) { + std::cerr << "MPI_THREAD_MULTIPLE support disabled.\n"; + } + } + if (comm == MPI_COMM_NULL) comm = MPI_COMM_WORLD; + + NET_CHECK(MPI_Comm_dup(comm, &thread_comm)); + NET_CHECK(MPI_Comm_size(thread_comm, &num_pes)); + NET_CHECK(MPI_Comm_rank(thread_comm, &my_pe)); +} + +void IPCBackend::team_destroy(roc_shmem_team_t team) { + IPCTeam *team_obj = get_internal_ipc_team(team); + + /* Mark the pool as available */ + int bit = team_obj->pool_index_; + int byte_i = bit / CHAR_BIT; + pool_bitmask_[byte_i] |= 1 << (bit % CHAR_BIT); + + team_obj->~IPCTeam(); + CHECK_HIP(hipFree(team_obj)); +} + +void IPCBackend::create_new_team([[maybe_unused]] Team *parent_team, + TeamInfo *team_info_wrt_parent, + TeamInfo *team_info_wrt_world, int num_pes, + int my_pe_in_new_team, MPI_Comm team_comm, + roc_shmem_team_t *new_team) { + /** + * Read the bit mask and find out a common index into + * the pool of available work arrays. + */ + NET_CHECK(MPI_Allreduce(pool_bitmask_, reduced_bitmask_, bitmask_size_, + MPI_CHAR, MPI_BAND, team_comm)); + + /* Pick the least significant non-zero bit (logical layout) in the reduced + * bitmask */ + auto max_num_teams{team_tracker.get_max_num_teams()}; + int common_index = get_ls_non_zero_bit(reduced_bitmask_, max_num_teams); + if (common_index < 0) { + /* No team available */ + abort(); + } + + /* Mark the team as taken (by unsetting the bit in the pool bitmask) */ + int byte = common_index / CHAR_BIT; + pool_bitmask_[byte] &= ~(1 << (common_index % CHAR_BIT)); + + /** + * Allocate device-side memory for team_world and + * construct a GPU_IB team in it + */ + GPUIBTeam *new_team_obj; + CHECK_HIP(hipMalloc(&new_team_obj, sizeof(IPCTeam))); + new (new_team_obj) + IPCTeam(this, team_info_wrt_parent, team_info_wrt_world, num_pes, + my_pe_in_new_team, team_comm, common_index); + + *new_team = get_external_team(new_team_obj); +} + +void IPCBackend::ctx_create(int64_t options, void **ctx) { + IPCHostContext *new_ctx{nullptr}; + new_ctx = new IPCHostContext(this, options); + *ctx = new_ctx; +} + +IPCHostContext *get_internal_ipc_net_ctx(Context *ctx) { + return reinterpret_cast(ctx); +} + +void IPCBackend::ctx_destroy(Context *ctx) { + IPCHostContext *ro_net_host_ctx{get_internal_ipc_net_ctx(ctx)}; + delete ro_net_host_ctx; +} + +void IPCBackend::reset_backend_stats() { + assert(false); +} + +void IPCBackend::dump_backend_stats() { + assert(false); +} + +void IPCBackend::initIPC() { + const auto &heap_bases{heap.get_heap_bases()}; + + ipcImpl.ipcHostInit(my_pe, heap_bases, + thread_comm); +} + +void IPCBackend::global_exit(int status) { + assert(false); +} + +void IPCBackend::teams_destroy() { + roc_shmem_free(barrier_pSync_pool); + roc_shmem_free(reduce_pSync_pool); + roc_shmem_free(bcast_pSync_pool); + roc_shmem_free(alltoall_pSync_pool); + roc_shmem_free(pWrk_pool); + roc_shmem_free(pAta_pool); + + free(pool_bitmask_); + free(reduced_bitmask_); +} + +void IPCBackend::roc_shmem_collective_init() { + /* + * Allocate heap space for barrier_sync + */ + size_t one_sync_size_bytes{sizeof(*barrier_sync)}; + size_t sync_size_bytes{one_sync_size_bytes * ROC_SHMEM_BARRIER_SYNC_SIZE}; + heap.malloc(reinterpret_cast(&barrier_sync), sync_size_bytes); + + /* + * Initialize the barrier synchronization array with default values. + */ + for (int i = 0; i < num_pes; i++) { + barrier_sync[i] = ROC_SHMEM_SYNC_VALUE; + } + + /* + * Make sure that all processing elements have done this before + * continuing. + */ + NET_CHECK(MPI_Barrier(thread_comm)); +} + +void IPCBackend::teams_init() { + /** + * Allocate pools for the teams sync and work arrary from the SHEAP. + */ + auto max_num_teams{team_tracker.get_max_num_teams()}; + barrier_pSync_pool = reinterpret_cast(roc_shmem_malloc( + sizeof(long) * ROC_SHMEM_BARRIER_SYNC_SIZE * max_num_teams)); + reduce_pSync_pool = reinterpret_cast(roc_shmem_malloc( + sizeof(long) * ROC_SHMEM_REDUCE_SYNC_SIZE * max_num_teams)); + bcast_pSync_pool = reinterpret_cast(roc_shmem_malloc( + sizeof(long) * ROC_SHMEM_BCAST_SYNC_SIZE * max_num_teams)); + alltoall_pSync_pool = reinterpret_cast(roc_shmem_malloc( + sizeof(long) * ROC_SHMEM_ALLTOALL_SYNC_SIZE * max_num_teams)); + + /* Accommodating for largest possible data type for pWrk */ + pWrk_pool = roc_shmem_malloc( + sizeof(double) * ROC_SHMEM_REDUCE_MIN_WRKDATA_SIZE * max_num_teams); + pAta_pool = roc_shmem_malloc(sizeof(double) * ROC_SHMEM_ATA_MAX_WRKDATA_SIZE * + max_num_teams); + + /** + * Initialize the sync arrays in the pool with default values. + */ + long *barrier_pSync, *reduce_pSync, *bcast_pSync, *alltoall_pSync; + for (int team_i = 0; team_i < max_num_teams; team_i++) { + barrier_pSync = reinterpret_cast( + &barrier_pSync_pool[team_i * ROC_SHMEM_BARRIER_SYNC_SIZE]); + reduce_pSync = reinterpret_cast( + &reduce_pSync_pool[team_i * ROC_SHMEM_REDUCE_SYNC_SIZE]); + bcast_pSync = reinterpret_cast( + &bcast_pSync_pool[team_i * ROC_SHMEM_BCAST_SYNC_SIZE]); + alltoall_pSync = reinterpret_cast( + &alltoall_pSync_pool[team_i * ROC_SHMEM_ALLTOALL_SYNC_SIZE]); + + for (int i = 0; i < ROC_SHMEM_BARRIER_SYNC_SIZE; i++) { + barrier_pSync[i] = ROC_SHMEM_SYNC_VALUE; + } + for (int i = 0; i < ROC_SHMEM_REDUCE_SYNC_SIZE; i++) { + reduce_pSync[i] = ROC_SHMEM_SYNC_VALUE; + } + for (int i = 0; i < ROC_SHMEM_BCAST_SYNC_SIZE; i++) { + bcast_pSync[i] = ROC_SHMEM_SYNC_VALUE; + } + for (int i = 0; i < ROC_SHMEM_ALLTOALL_SYNC_SIZE; i++) { + alltoall_pSync[i] = ROC_SHMEM_SYNC_VALUE; + } + } + + /** + * Initialize bit mask + * + * Logical: + * MSB..........................................................................LSB + * Physical: MSB...1st least significant 8 bits...LSB MSB...2nd least + * signifant 8 bits...LSB + * + * Description shows only a 2-byte long mask but idea extends to any + * arbitrary size. + */ + bitmask_size_ = (max_num_teams % CHAR_BIT) ? (max_num_teams / CHAR_BIT + 1) + : (max_num_teams / CHAR_BIT); + pool_bitmask_ = reinterpret_cast(malloc(bitmask_size_)); + reduced_bitmask_ = reinterpret_cast(malloc(bitmask_size_)); + + memset(pool_bitmask_, 0, bitmask_size_); + memset(reduced_bitmask_, 0, bitmask_size_); + /* Set all to available except the 0th one (reserved for TEAM_WORLD) */ + for (int bit_i = 1; bit_i < max_num_teams; bit_i++) { + int byte_i = bit_i / CHAR_BIT; + + pool_bitmask_[byte_i] |= 1 << (bit_i % CHAR_BIT); + } + + /** + * Make sure that all processing elements have done this before + * continuing. + */ + NET_CHECK(MPI_Barrier(thread_comm)); +} + +} // namespace rocshmem \ No newline at end of file diff --git a/src/ipc/backend_ipc.hpp b/src/ipc/backend_ipc.hpp new file mode 100644 index 0000000000..cc385f16ee --- /dev/null +++ b/src/ipc/backend_ipc.hpp @@ -0,0 +1,252 @@ +/****************************************************************************** + * 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_BACKEND_HPP_ +#define LIBRARY_SRC_IPC_BACKEND_HPP_ + +#include "../backend_bc.hpp" +#include "../containers/free_list_impl.hpp" +#include "../hdp_proxy.hpp" +#include "../memory/hip_allocator.hpp" +#include "ipc_backend_proxy.hpp" +#include "../context_incl.hpp" +#include "ipc_context_proxy.hpp" +#include "../ipc_policy.hpp" + +namespace rocshmem { + +class IPCBackend : public Backend { + const unsigned MAX_NUM_BLOCKS{65536}; + + public: + /** + * @copydoc Backend::Backend(unsigned) + */ + explicit IPCBackend(MPI_Comm comm); + + /** + * @copydoc Backend::~Backend() + */ + virtual ~IPCBackend(); + + __device__ bool create_ctx(int64_t options, roc_shmem_ctx_t *ctx); + + /** + * @brief Destroy a `roc_shmem_ctx_t` context and returns it back to the + * context free list. + */ + __device__ void destroy_ctx(roc_shmem_ctx_t *ctx); + + /** + * @copydoc Backend::ctx_create + */ + void ctx_create(int64_t options, void **ctx) override; + + /** + * @copydoc Backend::ctx_destroy + */ + void ctx_destroy(Context *ctx) override; + + /** + * @brief initialize MPI. + * + * IPC relies on MPI just to exchange the IPC_handle information. + * + * todo: remove the dependency on MPI and make it generic to PMI-X or just + * to OpenSHMEM to have support for both CPU and GPU + */ + void init_mpi_once(MPI_Comm comm); + + /** + * @brief Helper to initialize IPC interface. + */ + void initIPC(); + + /** + * @brief Allocation and initialization of backend contexts. + */ + void setup_ctxs(); + + /** + * @brief Abort the application. + * + * @param[in] status Exit code. + * + * @return void. + * + * @note This routine terminates the entire application. + */ + void global_exit(int status) override; + + /** + * @copydoc Backend::create_new_team + */ + void create_new_team(Team *parent_team, TeamInfo *team_info_wrt_parent, + TeamInfo *team_info_wrt_world, int num_pes, + int my_pe_in_new_team, MPI_Comm team_comm, + roc_shmem_team_t *new_team) override; + + /** + * @copydoc Backend::team_destroy(roc_shmem_team_t) + */ + void team_destroy(roc_shmem_team_t team) override; + + /** + * @brief Handle to device memory fields. + */ + IPCBackendProxyT ipc_backend_proxy{}; + + /** + * @brief The host-facing interface that will be used + * by all contexts of the ROBackend + */ + HostInterface *host_interface{nullptr}; + + /** + * @brief Scratchpad for the internal barrier algorithms. + */ + int64_t *barrier_sync{nullptr}; + + /** + * @brief Handle for raw memory for barrier sync + */ + long *barrier_pSync_pool{nullptr}; + + /** + * @brief Handle for raw memory for reduce sync + */ + long *reduce_pSync_pool{nullptr}; + + /** + * @brief Handle for raw memory for broadcast sync + */ + long *bcast_pSync_pool{nullptr}; + + /** + * @brief Handle for raw memory for alltoall sync + */ + long *alltoall_pSync_pool{nullptr}; + + /** + * @brief Handle for raw memory for work + */ + void *pWrk_pool{nullptr}; + + /** + * @brief Handle for raw memory for alltoall + */ + void *pAta_pool{nullptr}; + + protected: + /** + * @copydoc Backend::dump_backend_stats() + */ + void dump_backend_stats() override; + + /** + * @copydoc Backend::reset_backend_stats() + */ + void reset_backend_stats() override; + + /** + * @brief Allocates uncacheable host memory for the hdp policy. + * + * @note Internal data ownership is managed by the proxy + */ + HdpProxy hdp_proxy_{}; + + /** + * @brief Holds a copy of the default context for host functions + */ + std::unique_ptr default_host_ctx{nullptr}; + + /** + * @brief Allocate and initialize team world. + */ + void setup_team_world(); + + /** + * @brief Initialize the resources required to support teams + */ + void teams_init(); + + /** + * @brief Destruct the resources required to support teams + */ + void teams_destroy(); + + /** + * @brief Allocate and initialize barrier operation addresses on + * symmetric heap. + * + * When this method completes, the barrier_sync member will be available + * for use. + */ + void roc_shmem_collective_init(); + + private: + /** + * @brief Proxy for the default context + * + * @note Internal data ownership is managed by the proxy + */ + IPCDefaultContextProxyT default_context_proxy_; // init handled in constructor + + /** + * @brief An array of @ref ROContexts that backs the context FreeList. + */ + IPCContext *ctx_array{nullptr}; + + /** + * @brief A free-list containing contexts. + */ + FreeListProxy ctx_free_list{}; + + /** + * @brief Holds maximum number of contexts used in library + */ + size_t maximum_num_contexts_{1024}; + + /** + * @brief The bitmask representing the availability of teams in the pool + */ + char *pool_bitmask_{nullptr}; + + /** + * @brief Bitmask to store the reduced result of bitmasks on pariticipating + * PEs + * + * With no thread-safety for this bitmask, multithreaded creation of teams is + * not supported. + */ + char *reduced_bitmask_{nullptr}; + + /** + * @brief Size of the bitmask + */ + int bitmask_size_{-1}; + +}; + +} // namespace rocshmem + +#endif // LIBRARY_SRC_IPC_BACKEND_HPP_ \ No newline at end of file diff --git a/src/ipc/context_ipc_device.cpp b/src/ipc/context_ipc_device.cpp new file mode 100644 index 0000000000..6fbb362d73 --- /dev/null +++ b/src/ipc/context_ipc_device.cpp @@ -0,0 +1,156 @@ +/****************************************************************************** + * 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 "context_ipc_device.hpp" +#include "context_ipc_tmpl_device.hpp" + +#include +#include +#include + +#include +#include + +#include "config.h" // NOLINT(build/include_subdir) +#include "roc_shmem/roc_shmem.hpp" +#include "backend_ipc.hpp" + +namespace rocshmem { + +__host__ IPCContext::IPCContext(Backend *b) + : Context(b, false) { + IPCBackend *backend{static_cast(b)}; + ipcImpl_.ipc_bases = b->ipcImpl.ipc_bases; + ipcImpl_.shm_size = b->ipcImpl.shm_size; + + auto *bp{backend->ipc_backend_proxy.get()}; + + barrier_sync = backend->barrier_sync; + g_ret = bp->g_ret; + atomic_base_ptr = bp->atomic_ret->atomic_base_ptr; +} + +__device__ void IPCContext::threadfence_system() { +} + +__device__ void IPCContext::ctx_create() { +} + +__device__ void IPCContext::ctx_destroy(){ +} + +__device__ void IPCContext::putmem(void *dest, const void *source, size_t nelems, + int pe) { + uint64_t L_offset = + reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]; + ipcImpl_.ipcCopy(ipcImpl_.ipc_bases[pe] + L_offset, + const_cast(source), nelems); +} + +__device__ void IPCContext::getmem(void *dest, const void *source, size_t nelems, + int pe) { + const char *src_typed = reinterpret_cast(source); + uint64_t L_offset = + const_cast(src_typed) - ipcImpl_.ipc_bases[my_pe]; + ipcImpl_.ipcCopy(dest, ipcImpl_.ipc_bases[pe] + L_offset, nelems); +} + +__device__ void IPCContext::putmem_nbi(void *dest, const void *source, + size_t nelems, int pe) { + putmem(dest, source, nelems, pe); +} + +__device__ void IPCContext::getmem_nbi(void *dest, const void *source, + size_t nelems, int pe) { + getmem(dest, source, nelems, pe); +} + +__device__ void IPCContext::fence() { +} + +__device__ void IPCContext::fence(int pe) { +} + +__device__ void IPCContext::quiet() { +} + +__device__ void *IPCContext::shmem_ptr(const void *dest, int pe) { + void *ret = nullptr; + return ret; +} + +__device__ void IPCContext::putmem_wg(void *dest, const void *source, + size_t nelems, int pe) { + uint64_t L_offset = + reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]; + ipcImpl_.ipcCopy_wg(ipcImpl_.ipc_bases[pe] + L_offset, + const_cast(source), nelems); + __syncthreads(); +} + +__device__ void IPCContext::getmem_wg(void *dest, const void *source, + size_t nelems, int pe) { + const char *src_typed = reinterpret_cast(source); + uint64_t L_offset = + const_cast(src_typed) - ipcImpl_.ipc_bases[my_pe]; + ipcImpl_.ipcCopy_wg(dest, ipcImpl_.ipc_bases[pe] + L_offset, nelems); + __syncthreads(); +} + +__device__ void IPCContext::putmem_nbi_wg(void *dest, const void *source, + size_t nelems, int pe) { + putmem_wg(dest, source, nelems, pe); +} + +__device__ void IPCContext::getmem_nbi_wg(void *dest, const void *source, + size_t nelems, int pe) { + getmem_wg(dest, source, nelems, pe); +} + +__device__ void IPCContext::putmem_wave(void *dest, const void *source, + size_t nelems, int pe) { + uint64_t L_offset = + reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]; + ipcImpl_.ipcCopy_wave(ipcImpl_.ipc_bases[pe] + L_offset, + const_cast(source), nelems); +} + +__device__ void IPCContext::getmem_wave(void *dest, const void *source, + size_t nelems, int pe) { + const char *src_typed = reinterpret_cast(source); + uint64_t L_offset = + const_cast(src_typed) - ipcImpl_.ipc_bases[my_pe]; + ipcImpl_.ipcCopy_wave(dest, ipcImpl_.ipc_bases[pe] + L_offset, + nelems); +} + +__device__ void IPCContext::putmem_nbi_wave(void *dest, const void *source, + size_t nelems, int pe) { + putmem_wave(dest, source, nelems, pe); +} + +__device__ void IPCContext::getmem_nbi_wave(void *dest, const void *source, + size_t nelems, int pe) { + getmem_wave(dest, source, nelems, pe); +} + +} // namespace rocshmem diff --git a/src/ipc/context_ipc_device.hpp b/src/ipc/context_ipc_device.hpp new file mode 100644 index 0000000000..4cc6d08d56 --- /dev/null +++ b/src/ipc/context_ipc_device.hpp @@ -0,0 +1,239 @@ +/****************************************************************************** + * 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_CONTEXT_DEVICE_HPP_ +#define LIBRARY_SRC_IPC_CONTEXT_DEVICE_HPP_ + +#include "../context.hpp" + +namespace rocshmem { + +class IPCContext : public Context { + public: + __host__ IPCContext(Backend *b); + + __device__ IPCContext(Backend *b); + + __device__ void threadfence_system(); + + __device__ void ctx_create(); + + __device__ void ctx_destroy(); + + __device__ void putmem(void *dest, const void *source, size_t nelems, int pe); + + __device__ void getmem(void *dest, const void *source, size_t nelems, int pe); + + __device__ void putmem_nbi(void *dest, const void *source, size_t nelems, + int pe); + + __device__ void getmem_nbi(void *dest, const void *source, size_t size, + int pe); + + __device__ void fence(); + + __device__ void fence(int pe); + + __device__ void quiet(); + + __device__ void *shmem_ptr(const void *dest, int pe); + + __device__ void barrier_all(); + + __device__ void sync_all(); + + __device__ void sync(roc_shmem_team_t team); + + template + __device__ void p(T *dest, T value, int pe); + + template + __device__ void put(T *dest, const T *source, size_t nelems, int pe); + + template + __device__ void put_nbi(T *dest, const T *source, size_t nelems, int pe); + + template + __device__ T g(const T *source, int pe); + + template + __device__ void get(T *dest, const T *source, size_t nelems, int pe); + + template + __device__ void get_nbi(T *dest, const T *source, size_t nelems, int pe); + + // Atomic operations + template + __device__ void amo_add(void *dst, T value, int pe); + + template + __device__ void amo_set(void *dst, T value, int pe); + + template + __device__ T amo_swap(void *dst, T value, int pe); + + template + __device__ T amo_fetch_and(void *dst, T value, int pe); + + template + __device__ void amo_and(void *dst, T value, int pe); + + template + __device__ T amo_fetch_or(void *dst, T value, int pe); + + template + __device__ void amo_or(void *dst, T value, int pe); + + template + __device__ T amo_fetch_xor(void *dst, T value, int pe); + + template + __device__ void amo_xor(void *dst, T value, int pe); + + template + __device__ void amo_cas(void *dst, T value, T cond, int pe); + + template + __device__ T amo_fetch_add(void *dst, T value, int pe); + + template + __device__ T amo_fetch_cas(void *dst, T value, T cond, int pe); + + // Collectives + template + __device__ void to_all(T *dest, const T *source, int nreduce, int PE_start, + int logPE_stride, int PE_size, T *pWrk, + long *pSync); // NOLINT(runtime/int) + + template + __device__ void to_all(roc_shmem_team_t team, T *dest, const T *source, + int nreduce); + + template + __device__ void broadcast(roc_shmem_team_t team, T *dest, const T *source, + int nelems, int pe_root); + + template + __device__ void broadcast(T *dest, const T *source, int nelems, int pe_root, + int pe_start, int log_pe_stride, int pe_size, + long *p_sync); // NOLINT(runtime/int) + template + __device__ void alltoall(roc_shmem_team_t team, T *dest, const T *source, + int nelems); + template + __device__ void fcollect(roc_shmem_team_t team, T *dest, const T *source, + int nelems); + + + // Block/wave functions + __device__ void putmem_wg(void *dest, const void *source, size_t nelems, + int pe); + + __device__ void getmem_wg(void *dest, const void *source, size_t nelems, + int pe); + + __device__ void putmem_nbi_wg(void *dest, const void *source, size_t nelems, + int pe); + + __device__ void getmem_nbi_wg(void *dest, const void *source, size_t size, + int pe); + + __device__ void putmem_wave(void *dest, const void *source, size_t nelems, + int pe); + + __device__ void getmem_wave(void *dest, const void *source, size_t nelems, + int pe); + + __device__ void putmem_nbi_wave(void *dest, const void *source, size_t nelems, + int pe); + + __device__ void getmem_nbi_wave(void *dest, const void *source, size_t size, + int pe); + + template + __device__ void put_wg(T *dest, const T *source, size_t nelems, int pe); + + template + __device__ void put_nbi_wg(T *dest, const T *source, size_t nelems, int pe); + + template + __device__ void put_wave(T *dest, const T *source, size_t nelems, int pe); + + template + __device__ void put_nbi_wave(T *dest, const T *source, size_t nelems, int pe); + + template + __device__ void get_wg(T *dest, const T *source, size_t nelems, int pe); + + template + __device__ void get_nbi_wg(T *dest, const T *source, size_t nelems, int pe); + + + template + __device__ void get_wave(T *dest, const T *source, size_t nelems, int pe); + + template + __device__ void get_nbi_wave(T *dest, const T *source, size_t nelems, int pe); + + private: + + //context class has IpcImpl object (ipcImpl_) + IpcImpl *ipcImpl{nullptr}; + + uint64_t* atomic_base_ptr{nullptr}; + + char* g_ret; + + //internal functions used by collective operations + template + __device__ void internal_put_broadcast(T *dst, const T *src, int nelems, + int pe_root, int PE_start, + int logPE_stride, int PE_size); // NOLINT(runtime/int) + + template + __device__ void internal_get_broadcast(T *dst, const T *src, int nelems, + int pe_root); // NOLINT(runtime/int) + + template + __device__ void fcollect_linear(roc_shmem_team_t team, T *dest, + const T *source, int nelems); + + template + __device__ void alltoall_linear(roc_shmem_team_t team, T *dest, + const T *source, int nelems); + + __device__ void internal_sync(int pe, int PE_start, int stride, int PE_size, + int64_t *pSync); + + __device__ void internal_direct_barrier(int pe, int PE_start, int stride, + int n_pes, int64_t *pSync); + + __device__ void internal_atomic_barrier(int pe, int PE_start, int stride, + int n_pes, int64_t *pSync); + + //Temporary scratchpad memory used by internal barrier algorithms. + int64_t *barrier_sync{nullptr}; +}; + +} // namespace rocshmem + +#endif // LIBRARY_SRC_GPU_IB_CONTEXT_IB_DEVICE_HPP_ diff --git a/src/ipc/context_ipc_device_coll.cpp b/src/ipc/context_ipc_device_coll.cpp new file mode 100644 index 0000000000..804be2efbf --- /dev/null +++ b/src/ipc/context_ipc_device_coll.cpp @@ -0,0 +1,120 @@ +/****************************************************************************** + * 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 "roc_shmem/roc_shmem.hpp" +#include "../context_incl.hpp" +#include "context_ipc_tmpl_device.hpp" +#include "../util.hpp" +#include "ipc_team.hpp" + +namespace rocshmem { + +__device__ void IPCContext::internal_direct_barrier(int pe, int PE_start, + int stride, int n_pes, + int64_t *pSync) { + int64_t flag_val = 1; + if (pe == PE_start) { + // Go through all PE offsets (except current offset = 0) + // and wait until they all reach + for (size_t i = 1; i < n_pes; i++) { + wait_until(&pSync[i], ROC_SHMEM_CMP_EQ, flag_val); + pSync[i] = ROC_SHMEM_SYNC_VALUE; + } + threadfence_system(); + + // Announce to other PEs that all have reached + for (size_t i = 1, j = PE_start + stride; i < n_pes; ++i, j += stride) { + put_nbi(&pSync[0], &flag_val, 1, j); + } + } else { + // Mark current PE offset as reached + size_t pe_offset = (pe - PE_start) / stride; + put_nbi(&pSync[pe_offset], &flag_val, 1, PE_start); + wait_until(&pSync[0], ROC_SHMEM_CMP_EQ, flag_val); + pSync[0] = ROC_SHMEM_SYNC_VALUE; + threadfence_system(); + } +} + +__device__ void IPCContext::internal_atomic_barrier(int pe, int PE_start, + int stride, int n_pes, + int64_t *pSync) { + int64_t flag_val = 1; + if (pe == PE_start) { + wait_until(&pSync[0], ROC_SHMEM_CMP_EQ, (int64_t)(n_pes - 1)); + pSync[0] = ROC_SHMEM_SYNC_VALUE; + threadfence_system(); + + for (size_t i = 1, j = PE_start + stride; i < n_pes; ++i, j += stride) { + put_nbi(&pSync[0], &flag_val, 1, j); + } + } else { + amo_add(&pSync[0], flag_val, PE_start); + wait_until(&pSync[0], ROC_SHMEM_CMP_EQ, flag_val); + pSync[0] = ROC_SHMEM_SYNC_VALUE; + threadfence_system(); + } +} + +// Uses PE values that are relative to world +__device__ void IPCContext::internal_sync(int pe, int PE_start, int stride, + int PE_size, int64_t *pSync) { + __syncthreads(); + if (is_thread_zero_in_block()) { + if (PE_size < 64) { + internal_direct_barrier(pe, PE_start, stride, PE_size, pSync); + } else { + internal_atomic_barrier(pe, PE_start, stride, PE_size, pSync); + } + } + __threadfence(); + __syncthreads(); +} + +__device__ void IPCContext::sync(roc_shmem_team_t team) { + IPCTeam *team_obj = reinterpret_cast(team); + + /** + * Ensure that the stride is a multiple of 2. + */ + int log_pe_stride = static_cast(team_obj->tinfo_wrt_world->log_stride); + int pe = team_obj->my_pe_in_world; + int pe_start = team_obj->tinfo_wrt_world->pe_start; + int pe_stride = (1 << log_pe_stride); + int pe_size = team_obj->num_pes; + + internal_sync(pe, pe_start, pe_stride, pe_size, barrier_sync); +} + +__device__ void IPCContext::sync_all() { + internal_sync(my_pe, 0, 1, num_pes, barrier_sync); +} + +__device__ void IPCContext::barrier_all() { + if (is_thread_zero_in_block()) { + quiet(); + } + sync_all(); + __syncthreads(); +} + +} // namespace rocshmem diff --git a/src/ipc/context_ipc_host.cpp b/src/ipc/context_ipc_host.cpp new file mode 100644 index 0000000000..0d3464f33d --- /dev/null +++ b/src/ipc/context_ipc_host.cpp @@ -0,0 +1,85 @@ +/****************************************************************************** + * 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 "context_ipc_host.hpp" + +#include + +#include "config.h" // NOLINT(build/include_subdir) +#include "../backend_type.hpp" +#include "../context_incl.hpp" +#include "backend_ipc.hpp" +#include "../host/host.hpp" + +namespace rocshmem { + +__host__ IPCHostContext::IPCHostContext(Backend *backend, + [[maybe_unused]] int64_t options) + : Context(backend, true) { + IPCBackend *b{static_cast(backend)}; + + host_interface = b->host_interface; + + context_window_info = host_interface->acquire_window_context(); +} + +__host__ IPCHostContext::~IPCHostContext() { + host_interface->release_window_context(context_window_info); +} + +__host__ void IPCHostContext::putmem_nbi(void *dest, const void *source, + size_t nelems, int pe) { + host_interface->putmem_nbi(dest, source, nelems, pe, context_window_info); +} + +__host__ void IPCHostContext::getmem_nbi(void *dest, const void *source, + size_t nelems, int pe) { + host_interface->getmem_nbi(dest, source, nelems, pe, context_window_info); +} + +__host__ void IPCHostContext::putmem(void *dest, const void *source, + size_t nelems, int pe) { + host_interface->putmem(dest, source, nelems, pe, context_window_info); +} + +__host__ void IPCHostContext::getmem(void *dest, const void *source, + size_t nelems, int pe) { + host_interface->getmem(dest, source, nelems, pe, context_window_info); +} + +__host__ void IPCHostContext::fence() { + host_interface->fence(context_window_info); +} + +__host__ void IPCHostContext::quiet() { + host_interface->quiet(context_window_info); +} + +__host__ void IPCHostContext::sync_all() { + host_interface->sync_all(context_window_info); +} + +__host__ void IPCHostContext::barrier_all() { + host_interface->barrier_all(context_window_info); +} + +} // namespace rocshmem diff --git a/src/ipc/context_ipc_host.hpp b/src/ipc/context_ipc_host.hpp new file mode 100644 index 0000000000..f9421ca210 --- /dev/null +++ b/src/ipc/context_ipc_host.hpp @@ -0,0 +1,149 @@ +/****************************************************************************** + * 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_CONTEXT_HOST_HPP_ +#define LIBRARY_SRC_IPC_CONTEXT_HOST_HPP_ + +#include "../context.hpp" + +namespace rocshmem { + +class IPCHostContext : public Context { + public: + __host__ IPCHostContext(Backend *b, int64_t options); + + __host__ ~IPCHostContext(); + + template + __host__ void p(T *dest, T value, int pe); + + template + __host__ T g(const T *source, int pe); + + template + __host__ void put(T *dest, const T *source, size_t nelems, int pe); + + template + __host__ void get(T *dest, const T *source, size_t nelems, int pe); + + template + __host__ void put_nbi(T *dest, const T *source, size_t nelems, int pe); + + template + __host__ void get_nbi(T *dest, const T *source, size_t nelems, int pe); + + __host__ void putmem(void *dest, const void *source, size_t nelems, int pe); + + __host__ void getmem(void *dest, const void *source, size_t nelems, int pe); + + __host__ void putmem_nbi(void *dest, const void *source, size_t nelems, + int pe); + + __host__ void getmem_nbi(void *dest, const void *source, size_t size, int pe); + + template + __host__ void amo_add(void *dst, T value, int pe); + + template + __host__ void amo_cas(void *dst, T value, T cond, int pe); + + template + __host__ T amo_fetch_add(void *dst, T value, int pe); + + template + __host__ T amo_fetch_cas(void *dst, T value, T cond, int pe); + + __host__ void fence(); + + __host__ void quiet(); + + __host__ void barrier_all(); + + __host__ void sync_all(); + + template + __host__ void broadcast(T *dest, const T *source, int nelems, int pe_root, + int pe_start, int log_pe_stride, int pe_size, + long *p_sync); + + template + __host__ void broadcast(roc_shmem_team_t team, T *dest, const T *source, + int nelems, int pe_root); + + template + __host__ void to_all(T *dest, const T *source, int nreduce, int pe_start, + int log_pe_stride, int pe_size, T *p_wrk, + long *p_sync); + + template + __host__ void to_all(roc_shmem_team_t team, T *dest, const T *source, + int nreduce); + + template + __host__ void wait_until(T *ptr, roc_shmem_cmps cmp, T val); + + template + __host__ size_t wait_until_any(T* ptr, size_t nelems, + const int *status, + roc_shmem_cmps cmp, T val); + + template + __host__ void wait_until_all(T* ptr, size_t nelems, + const int *status, + roc_shmem_cmps cmp, T val); + + template + __host__ size_t wait_until_some(T* ptr, size_t nelems, + size_t* indices, + const int *status, + roc_shmem_cmps cmp, T val); + + template + __host__ void wait_until_all_vector(T* ptr, size_t nelems, + const int *status, + roc_shmem_cmps cmp, T* vals); + + template + __host__ size_t wait_until_any_vector(T* ptr, size_t nelems, + const int *status, + roc_shmem_cmps cmp, T* vals); + + template + __host__ size_t wait_until_some_vector(T* ptr, size_t nelems, + size_t* indices, + const int *status, + roc_shmem_cmps cmp, T* vals); + + template + __host__ int test(T *ptr, roc_shmem_cmps cmp, T val); + + public: + /* Pointer to the backend's host interface */ + HostInterface *host_interface{nullptr}; + + /* An MPI Window implements a context */ + WindowInfo *context_window_info{nullptr}; +}; + +} // namespace rocshmem + +#endif // LIBRARY_SRC_IPC_CONTEXT_HOST_HPP_ diff --git a/src/ipc/context_ipc_tmpl_device.hpp b/src/ipc/context_ipc_tmpl_device.hpp new file mode 100644 index 0000000000..91bdbd45e7 --- /dev/null +++ b/src/ipc/context_ipc_tmpl_device.hpp @@ -0,0 +1,345 @@ +/****************************************************************************** + * 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_CONTEXT_TMPL_DEVICE_HPP_ +#define LIBRARY_SRC_IPC_CONTEXT_TMPL_DEVICE_HPP_ + +#include "config.h" // NOLINT(build/include_subdir) +#include "roc_shmem/roc_shmem.hpp" +#include "context_ipc_device.hpp" +#include "../util.hpp" +#include "ipc_team.hpp" + +namespace rocshmem { + +/****************************************************************************** + ************************** TEMPLATE SPECIALIZATIONS ************************** + *****************************************************************************/ +template +__device__ void IPCContext::p(T *dest, T value, int pe) { + putmem_nbi(dest, &value, sizeof(T), pe); +} + +template +__device__ void IPCContext::put(T *dest, const T *source, size_t nelems, + int pe) { + putmem(dest, source, nelems * sizeof(T), pe); +} + +template +__device__ void IPCContext::put_nbi(T *dest, const T *source, size_t nelems, + int pe) { + putmem_nbi(dest, source, sizeof(T) * nelems, pe); +} + +template +__device__ T IPCContext::g(const T *source, int pe) { + T ret; + return ret; +} + +template +__device__ void IPCContext::get(T *dest, const T *source, size_t nelems, + int pe) { + getmem(dest, source, sizeof(T) * nelems, pe); +} + +template +__device__ void IPCContext::get_nbi(T *dest, const T *source, size_t nelems, + int pe) { + getmem_nbi(dest, source, sizeof(T) * nelems, pe); +} + +// Atomics +template +__device__ void IPCContext::amo_add(void *dest, T value, int pe) { + uint64_t L_offset = + reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]; + ipcImpl_.ipcAMOAdd( + reinterpret_cast(ipcImpl_.ipc_bases[pe] + L_offset), value); +} + +template +__device__ void IPCContext::amo_set(void *dest, T value, int pe) { + uint64_t L_offset = + reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]; + ipcImpl_.ipcAMOSet( + reinterpret_cast(ipcImpl_.ipc_bases[pe] + L_offset), value); +} + +template +__device__ T IPCContext::amo_swap(void *dst, T value, int pe) { + assert(false); + return 0; +} + +template +__device__ T IPCContext::amo_fetch_and(void *dst, T value, int pe) { + assert(false); + return 0; +} + +template +__device__ void IPCContext::amo_and(void *dst, T value, int pe) { + assert(false); +} + +template +__device__ T IPCContext::amo_fetch_or(void *dst, T value, int pe) { + assert(false); + return 0; +} + +template +__device__ void IPCContext::amo_or(void *dst, T value, int pe) { + assert(false); +} + +template +__device__ T IPCContext::amo_fetch_xor(void *dst, T value, int pe) { + assert(false); + return 0; +} + +template +__device__ void IPCContext::amo_xor(void *dst, T value, int pe) { + assert(false); +} + +template +__device__ void IPCContext::amo_cas(void *dest, T value, T cond, int pe) { + uint64_t L_offset = + reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]; + ipcImpl_.ipcAMOCas( + reinterpret_cast(ipcImpl_.ipc_bases[pe] + L_offset), cond, + value); +} + +template +__device__ T IPCContext::amo_fetch_add(void *dest, T value, int pe) { + uint64_t L_offset = + reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]; + return ipcImpl_.ipcAMOFetchAdd( + reinterpret_cast(ipcImpl_.ipc_bases[pe] + L_offset), value); +} + +template +__device__ T IPCContext::amo_fetch_cas(void *dest, T value, T cond, int pe) { + uint64_t L_offset = + reinterpret_cast(dest) - ipcImpl_.ipc_bases[my_pe]; + return ipcImpl_.ipcAMOFetchCas( + reinterpret_cast(ipcImpl_.ipc_bases[pe] + L_offset), cond, + value); +} + +// Collectives +template +__device__ void IPCContext::to_all(roc_shmem_team_t team, T *dest, + const T *source, int nreduce) { + //to_all(dest, source, nreduce, pe_start, log_pe_stride, pe_size, pWrk, + // p_sync); +} + +template +__device__ void IPCContext::to_all(T *dest, const T *source, int nreduce, + int PE_start, int logPE_stride, + int PE_size, T *pWrk, + long *pSync) { // NOLINT(runtime/int) +} + +template +__device__ void IPCContext::internal_put_broadcast( + T *dst, const T *src, int nelems, int pe_root, int pe_start, + int log_pe_stride, int pe_size) { // NOLINT(runtime/int) + if (my_pe == pe_root) { + int stride = 1 << log_pe_stride; + int finish = pe_start + stride * pe_size; + for (int i = pe_start; i < finish; i += stride) { + if (i != my_pe) { + put_nbi_wg(dst, src, nelems, i); + } + } + } +} + +template +__device__ void IPCContext::internal_get_broadcast( + T *dst, const T *src, int nelems, int pe_root) { // NOLINT(runtime/int) + if (my_pe != pe_root) { + get_wg(dst, src, nelems, pe_root); + } +} + +template +__device__ void IPCContext::broadcast(roc_shmem_team_t team, T *dst, + const T *src, int nelems, int pe_root) { + IPCTeam *team_obj = reinterpret_cast(team); + + /** + * Ensure that the stride is a multiple of 2 . + */ + int log_pe_stride = static_cast(team_obj->tinfo_wrt_world->log_stride); + int pe_start = team_obj->tinfo_wrt_world->pe_start; + int pe_size = team_obj->tinfo_wrt_world->size; + long *p_sync = team_obj->bcast_pSync; + + // Passed pe_root is relative to team, convert to world root + int pe_root_world = team_obj->get_pe_in_world(pe_root); + broadcast(dst, src, nelems, pe_root_world, pe_start, log_pe_stride, + pe_size, p_sync); +} + +template +__device__ void IPCContext::broadcast(T *dst, const T *src, int nelems, + int pe_root, int pe_start, + int log_pe_stride, int pe_size, + long *p_sync) { // NOLINT(runtime/int) + if (num_pes < 4) { + internal_put_broadcast(dst, src, nelems, pe_root, pe_start, log_pe_stride, + pe_size); + } else { + internal_get_broadcast(dst, src, nelems, pe_root); + } + + // Synchronize on completion of broadcast + internal_sync(my_pe, pe_start, (1 << log_pe_stride), pe_size, p_sync); +} + +template +__device__ void IPCContext::alltoall(roc_shmem_team_t team, T *dst, + const T *src, int nelems) { + alltoall_linear(team, dst, src, nelems); +} + +template +__device__ void IPCContext::alltoall_linear(roc_shmem_team_t team, T *dst, + const T *src, int nelems) { + IPCTeam *team_obj = reinterpret_cast(team); + + /** + * Ensure that the stride is a multiple of 2 + */ + int log_pe_stride = static_cast(team_obj->tinfo_wrt_world->log_stride); + int pe_start = team_obj->tinfo_wrt_world->pe_start; + int pe_size = team_obj->num_pes; + int stride = 1 << log_pe_stride; + long *pSync = team_obj->alltoall_pSync; + int my_pe_in_team = team_obj->my_pe; + + // Have each PE put their designated data to the other PEs + for (int j = 0; j < pe_size; j++) { + int dest_pe = team_obj->get_pe_in_world(j); + put_nbi_wg(&dst[my_pe_in_team * nelems], &src[j * nelems], nelems, dest_pe); + } + if (is_thread_zero_in_block()) { + quiet(); + } + // wait until everyone has obtained their designated data + internal_sync(my_pe, pe_start, stride, pe_size, pSync); +} + +template +__device__ void IPCContext::fcollect(roc_shmem_team_t team, T *dst, + const T *src, int nelems) { + fcollect_linear(team, dst, src, nelems); +} + +template +__device__ void IPCContext::fcollect_linear(roc_shmem_team_t team, T *dst, + const T *src, int nelems) { + IPCTeam *team_obj = reinterpret_cast(team); + + /** + * Ensure that the stride is a multiple of 2. + */ + int log_pe_stride = static_cast(team_obj->tinfo_wrt_world->log_stride); + int pe_start = team_obj->tinfo_wrt_world->pe_start; + int pe_size = team_obj->num_pes; + int stride = 1 << log_pe_stride; + long *pSync = team_obj->alltoall_pSync; + int my_pe_in_team = team_obj->my_pe; + + // Have each PE put their designated data to the other PEs + for (int j = 0; j < pe_size; j++) { + int dest_pe = team_obj->get_pe_in_world(j); + put_nbi_wg(&dst[my_pe_in_team * nelems], src, nelems, dest_pe); + } + + if (is_thread_zero_in_block()) { + quiet(); + } + // wait until everyone has obtained their designated data + internal_sync(my_pe, pe_start, stride, pe_size, pSync); +} + +// Block/wave functions +template +__device__ void IPCContext::put_wg(T *dest, const T *source, size_t nelems, + int pe) { + putmem_wg(dest, source, nelems * sizeof(T), pe); +} + +template +__device__ void IPCContext::put_nbi_wg(T *dest, const T *source, + size_t nelems, int pe) { + putmem_nbi_wg(dest, source, nelems * sizeof(T), pe); +} + + template +__device__ void IPCContext::put_wave(T *dest, const T *source, size_t nelems, + int pe) { + putmem_wave(dest, source, nelems * sizeof(T), pe); +} + +template +__device__ void IPCContext::put_nbi_wave(T *dest, const T *source, + size_t nelems, int pe) { + putmem_nbi_wave(dest, source, nelems * sizeof(T), pe); +} + +template +__device__ void IPCContext::get_wg(T *dest, const T *source, size_t nelems, + int pe) { + getmem_wg(dest, source, nelems * sizeof(T), pe); +} + +template +__device__ void IPCContext::get_nbi_wg(T *dest, const T *source, + size_t nelems, int pe) { + getmem_nbi_wg(dest, source, nelems * sizeof(T), pe); +} + +template +__device__ void IPCContext::get_wave(T *dest, const T *source, size_t nelems, + int pe) { + getmem_wave(dest, source, nelems * sizeof(T), pe); +} + +template +__device__ void IPCContext::get_nbi_wave(T *dest, const T *source, + size_t nelems, int pe) { + getmem_nbi_wave(dest, source, nelems * sizeof(T), pe); +} + +} // namespace rocshmem + +#endif // LIBRARY_SRC_IPC_CONTEXT_TMPL_DEVICE_HPP_ diff --git a/src/ipc/context_ipc_tmpl_host.hpp b/src/ipc/context_ipc_tmpl_host.hpp new file mode 100644 index 0000000000..76484de34d --- /dev/null +++ b/src/ipc/context_ipc_tmpl_host.hpp @@ -0,0 +1,173 @@ +/****************************************************************************** + * 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_CONTEXT_TMPL_HOST_HPP_ +#define LIBRARY_SRC_IPC_CONTEXT_TMPL_HOST_HPP_ + +#include "config.h" // NOLINT(build/include_subdir) +#include "../host/host_templates.hpp" + +namespace rocshmem { + +template +__host__ void IPCHostContext::p(T *dest, T value, int pe) { + host_interface->p(dest, value, pe, context_window_info); +} + +template +__host__ T IPCHostContext::g(const T *source, int pe) { + return host_interface->g(source, pe, context_window_info); +} + +template +__host__ void IPCHostContext::put(T *dest, const T *source, size_t nelems, + int pe) { + host_interface->put(dest, source, nelems, pe, context_window_info); +} + +template +__host__ void IPCHostContext::get(T *dest, const T *source, size_t nelems, + int pe) { + host_interface->get(dest, source, nelems, pe, context_window_info); +} + +template +__host__ void IPCHostContext::put_nbi(T *dest, const T *source, size_t nelems, + int pe) { + host_interface->put_nbi(dest, source, nelems, pe, context_window_info); +} + +template +__host__ void IPCHostContext::get_nbi(T *dest, const T *source, size_t nelems, + int pe) { + host_interface->get_nbi(dest, source, nelems, pe, context_window_info); +} + +template +__host__ void IPCHostContext::amo_add(void *dst, T value, int pe) { + host_interface->amo_add(dst, value, pe, context_window_info); +} + +template +__host__ void IPCHostContext::amo_cas(void *dst, T value, T cond, int pe) { + host_interface->amo_cas(dst, value, cond, pe, context_window_info); +} + +template +__host__ T IPCHostContext::amo_fetch_add(void *dst, T value, int pe) { + return host_interface->amo_fetch_add(dst, value, pe, context_window_info); +} + +template +__host__ T IPCHostContext::amo_fetch_cas(void *dst, T value, T cond, int pe) { + return host_interface->amo_fetch_cas(dst, value, cond, pe, + context_window_info); +} + +template +__host__ void IPCHostContext::broadcast( + T *dest, const T *source, int nelems, int pe_root, int pe_start, + int log_pe_stride, int pe_size, + long *p_sync) { // NOLINT(runtime/int) + host_interface->broadcast(dest, source, nelems, pe_root, pe_start, + log_pe_stride, pe_size, p_sync); +} + +template +__host__ void IPCHostContext::broadcast(roc_shmem_team_t team, T *dest, + const T *source, int nelems, + int pe_root) { + host_interface->broadcast(team, dest, source, nelems, pe_root); +} + +template +__host__ void IPCHostContext::to_all(T *dest, const T *source, int nreduce, + int pe_start, int log_pe_stride, + int pe_size, T *p_wrk, + long *p_sync) { // NOLINT(runtime/int) + host_interface->to_all(dest, source, nreduce, pe_start, log_pe_stride, + pe_size, p_wrk, p_sync); +} + +template +__host__ void IPCHostContext::to_all(roc_shmem_team_t team, T *dest, + const T *source, int nreduce) { + host_interface->to_all(team, dest, source, nreduce); +} + +template +__host__ void IPCHostContext::wait_until(T *ptr, roc_shmem_cmps cmp, T val) { + host_interface->wait_until(ptr, cmp, val, context_window_info); +} + +template +__host__ void IPCHostContext::wait_until_all(T *ptr, size_t nelems, + const int* status, + roc_shmem_cmps cmp, T val) { + host_interface->wait_until_all(ptr, nelems, status, cmp, val, context_window_info); +} + +template +__host__ size_t IPCHostContext::wait_until_any(T *ptr, size_t nelems, + const int* status, + roc_shmem_cmps cmp, T val) { + return host_interface->wait_until_any(ptr, nelems, status, cmp, val, context_window_info); +} + +template +__host__ size_t IPCHostContext::wait_until_some(T *ptr, size_t nelems, + size_t* indices, + const int* status, + roc_shmem_cmps cmp, T val) { + return host_interface->wait_until_some(ptr, nelems, indices, status, cmp, val, context_window_info); +} + +template +__host__ void IPCHostContext::wait_until_all_vector(T *ptr, size_t nelems, + const int* status, + roc_shmem_cmps cmp, T* vals) { + host_interface->wait_until_all_vector(ptr, nelems, status, cmp, vals, context_window_info); +} + +template +__host__ size_t IPCHostContext::wait_until_any_vector(T *ptr, size_t nelems, + const int* status, + roc_shmem_cmps cmp, T* vals) { + return host_interface->wait_until_any_vector(ptr, nelems, status, cmp, vals, context_window_info); +} + +template +__host__ size_t IPCHostContext::wait_until_some_vector(T *ptr, size_t nelems, + size_t* indices, + const int* status, + roc_shmem_cmps cmp, T* vals) { + return host_interface->wait_until_some_vector(ptr, nelems, indices, status, cmp, vals, context_window_info); +} + +template +__host__ int IPCHostContext::test(T *ptr, roc_shmem_cmps cmp, T val) { + return host_interface->test(ptr, cmp, val, context_window_info); +} + +} // namespace rocshmem + +#endif // LIBRARY_SRC_IPC_CONTEXT_TMPL_HOST_HPP_ diff --git a/src/ipc/ipc_backend_proxy.hpp b/src/ipc/ipc_backend_proxy.hpp new file mode 100644 index 0000000000..0077fb053d --- /dev/null +++ b/src/ipc/ipc_backend_proxy.hpp @@ -0,0 +1,69 @@ +/****************************************************************************** + * 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_BACKEND_PROXY_HPP_ +#define LIBRARY_SRC_IPC_BACKEND_PROXY_HPP_ + +#include "../device_proxy.hpp" +#include "../atomic_return.hpp" + +namespace rocshmem { + +struct IPCBackendRegister { + char *g_ret{nullptr}; + atomic_ret_t *atomic_ret{nullptr}; + SymmetricHeap *heap_ptr{nullptr}; +}; + +template +class IPCBackendProxy { + using ProxyT = DeviceProxy; + + public: + /* + * Placement new the memory which is allocated by proxy_ + */ + IPCBackendProxy() { new (proxy_.get()) IPCBackendRegister(); } + + /* + * Since placement new is called in the constructor, then + * delete must be called manually. + */ + ~IPCBackendProxy() { proxy_.get()->~IPCBackendRegister(); } + + /* + * @brief Provide access to the memory referenced by the proxy + */ + __host__ __device__ IPCBackendRegister *get() { return proxy_.get(); } + + private: + /* + * @brief Memory managed by the lifetime of this object + */ + ProxyT proxy_{}; +}; + +using IPCBackendProxyT = IPCBackendProxy; + +} // namespace rocshmem + +#endif // #define LIBRARY_SRC_IPC_BACKEND_PROXY_HPP_ \ No newline at end of file diff --git a/src/ipc/ipc_context_proxy.hpp b/src/ipc/ipc_context_proxy.hpp new file mode 100644 index 0000000000..867f199094 --- /dev/null +++ b/src/ipc/ipc_context_proxy.hpp @@ -0,0 +1,90 @@ +/****************************************************************************** + * 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_CONTEXT_PROXY_HPP_ +#define LIBRARY_SRC_IPC_CONTEXT_PROXY_HPP_ + + +#include "../device_proxy.hpp" +#include "backend_ipc.hpp" + +namespace rocshmem { + +class IPCBackend; + +template +class IPCDefaultContextProxy { + using ProxyT = DeviceProxy; + + public: + IPCDefaultContextProxy() = default; + + /* + * Placement new the memory which is allocated by proxy_ + */ + explicit IPCDefaultContextProxy(IPCBackend* backend) : constructed_{true} { + auto ctx{proxy_.get()}; + new (ctx) IPCContext(reinterpret_cast(backend)); + roc_shmem_ctx_t local{ctx, nullptr}; + set_internal_ctx(&local); + } + + /* + * Since placement new is called in the constructor, then + * delete must be called manually. + */ + ~IPCDefaultContextProxy() { + if (constructed_) { + proxy_.get()->~IPCContext(); + } + } + + IPCDefaultContextProxy(const IPCDefaultContextProxy& other) = delete; + + IPCDefaultContextProxy& operator=(const IPCDefaultContextProxy& other) = delete; + + IPCDefaultContextProxy(IPCDefaultContextProxy&& other) = default; + + IPCDefaultContextProxy& operator=(IPCDefaultContextProxy&& other) = default; + + /* + * @brief Provide access to the memory referenced by the proxy + */ + __host__ __device__ Context* get() { return proxy_.get(); } + + private: + /* + * @brief Memory managed by the lifetime of this object + */ + ProxyT proxy_{}; + + /* + * @brief denotes if an objects was constructed in proxy + */ + bool constructed_{false}; +}; + +using IPCDefaultContextProxyT = IPCDefaultContextProxy; + +} // namespace rocshmem + +#endif // LIBRARY_SRC_IPC_CONTEXT_PROXY_HPP_ \ No newline at end of file diff --git a/src/ipc/ipc_team.cpp b/src/ipc/ipc_team.cpp new file mode 100644 index 0000000000..433ca32d4f --- /dev/null +++ b/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/src/ipc/ipc_team.hpp b/src/ipc/ipc_team.hpp new file mode 100644 index 0000000000..e931adcfe4 --- /dev/null +++ b/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/src/ipc_policy.cpp b/src/ipc_policy.cpp index 4bf483735d..2c3a96fa7c 100644 --- a/src/ipc_policy.cpp +++ b/src/ipc_policy.cpp @@ -50,7 +50,6 @@ __host__ void IpcOnImpl::ipcHostInit(int my_pe, const HEAP_BASES_T &heap_bases, /* * Figure out how this process' rank among local processes. */ - int shm_rank; MPI_Comm_rank(shmcomm, &shm_rank); /* @@ -92,7 +91,6 @@ __host__ void IpcOnImpl::ipcHostInit(int my_pe, const HEAP_BASES_T &heap_bases, void **ipc_base_uncast = reinterpret_cast(&ipc_base[i]); CHECK_HIP(hipIpcOpenMemHandle(ipc_base_uncast, vec_ipc_handle[i], hipIpcMemLazyEnablePeerAccess)); - // TODO(bpotter): add some error checking here if happens to fail } else { ipc_base[i] = base_heap; } @@ -110,6 +108,15 @@ __host__ void IpcOnImpl::ipcHostInit(int my_pe, const HEAP_BASES_T &heap_bases, free(vec_ipc_handle); } +__host__ void IpcOnImpl::ipcHostStop() { + for (size_t i = 0; i < shm_size; i++) { + if (i != shm_rank) { + CHECK_HIP(hipIpcCloseMemHandle(ipc_bases[i])); + } + } + CHECK_HIP(hipFree(ipc_bases)); +} + __device__ void IpcOnImpl::ipcCopy(void *dst, void *src, size_t size) { memcpy(dst, src, size); } diff --git a/src/ipc_policy.hpp b/src/ipc_policy.hpp index f6bb49d409..c0190198ca 100644 --- a/src/ipc_policy.hpp +++ b/src/ipc_policy.hpp @@ -42,6 +42,8 @@ class IpcOnImpl { using HEAP_BASES_T = std::vector>; public: + int shm_rank{0}; + uint32_t shm_size{0}; char **ipc_bases{nullptr}; @@ -49,6 +51,8 @@ class IpcOnImpl { __host__ void ipcHostInit(int my_pe, const HEAP_BASES_T &heap_bases, MPI_Comm thread_comm); + __host__ void ipcHostStop(); + __device__ bool isIpcAvailable(int my_pe, int target_pe) { return my_pe / shm_size == target_pe / shm_size; } @@ -115,6 +119,8 @@ class IpcOffImpl { __host__ void ipcHostInit(int my_pe, const HEAP_BASES_T &heap_bases, MPI_Comm thread_comm) {} + __host__ void ipcHostStop() {} + __device__ bool isIpcAvailable(int my_pe, int target_pe) { return false; } __device__ void ipcGpuInit(Backend *roc_shmem_handle, Context *ctx, diff --git a/src/roc_shmem.cpp b/src/roc_shmem.cpp index 8c61d6c864..32ef91bbda 100644 --- a/src/roc_shmem.cpp +++ b/src/roc_shmem.cpp @@ -39,12 +39,14 @@ #ifdef USE_GPU_IB #include "gpu_ib/backend_ib.hpp" #include "gpu_ib/context_ib_tmpl_host.hpp" -#endif -#include "mpi_init_singleton.hpp" -#ifndef USE_GPU_IB +#elif defined(USE_RO) #include "reverse_offload/backend_ro.hpp" #include "reverse_offload/context_ro_tmpl_host.hpp" +#else +#include "ipc/backend_ipc.hpp" +#include "ipc/context_ipc_tmpl_host.hpp" #endif +#include "mpi_init_singleton.hpp" #include "team.hpp" #include "templates_host.hpp" #include "util.hpp" @@ -82,12 +84,15 @@ roc_shmem_ctx_t ROC_SHMEM_HOST_CTX_DEFAULT; rocm_init(); -#ifndef USE_GPU_IB +#ifdef USE_GPU_IB + CHECK_HIP(hipHostMalloc(&backend, sizeof(GPUIBBackend))); + backend = new (backend) GPUIBBackend(comm); +#elif defined(USE_RO) CHECK_HIP(hipHostMalloc(&backend, sizeof(ROBackend))); backend = new (backend) ROBackend(comm); #else - CHECK_HIP(hipHostMalloc(&backend, sizeof(GPUIBBackend))); - backend = new (backend) GPUIBBackend(comm); + CHECK_HIP(hipHostMalloc(&backend, sizeof(IPCBackend))); + backend = new (backend) IPCBackend(comm); #endif if (!backend) { diff --git a/src/roc_shmem_gpu.cpp b/src/roc_shmem_gpu.cpp index 269a98a28a..086a6deed1 100644 --- a/src/roc_shmem_gpu.cpp +++ b/src/roc_shmem_gpu.cpp @@ -51,8 +51,10 @@ #ifdef USE_GPU_IB #include "gpu_ib/context_ib_tmpl_device.hpp" -#else +#elif defined(USE_RO) #include "reverse_offload/context_ro_tmpl_device.hpp" +#else +#include "ipc/context_ipc_tmpl_device.hpp" #endif /****************************************************************************** diff --git a/src/team.cpp b/src/team.cpp index fefb520328..21ed00037c 100644 --- a/src/team.cpp +++ b/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/src/team.hpp b/src/team.hpp index 56475e799b..0dfaac6ae3 100644 --- a/src/team.hpp +++ b/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); diff --git a/src/util.hpp b/src/util.hpp index c967ec1a9c..c209750d49 100644 --- a/src/util.hpp +++ b/src/util.hpp @@ -149,7 +149,7 @@ __device__ void gpu_dprintf(const char* fmt, const Args&... args) { while (atomicCAS(print_lock, 0, 1) == 1) { } - printf("WG (%lu, %lu, %lu) TH (%lu, %lu, %lu) ", hipBlockIdx_x, + printf("WG (%u, %u, %u) TH (%u, %u, %u) ", hipBlockIdx_x, hipBlockIdx_y, hipBlockIdx_z, hipThreadIdx_x, hipThreadIdx_y, hipThreadIdx_z); printf(fmt, args...); @@ -216,19 +216,34 @@ __device__ __forceinline__ void memcpy_wg(void* dst, void* src, size_t size) { } __device__ __forceinline__ void memcpy_wave(void* dst, void* src, size_t size) { - uint8_t* dst_bytes{static_cast(dst)}; - uint8_t* src_bytes{static_cast(src)}; + int wave_tid = get_flat_block_id() % WF_SIZE; + int wave_size{wave_SZ()}; int cpy_size{}; - int thread_id{get_flat_block_id()}; + uint8_t* dst_bytes{nullptr}; + uint8_t* dst_def{nullptr}; + uint8_t* src_bytes{nullptr}; + uint8_t* src_def{nullptr}; + + dst_def = reinterpret_cast(dst); + src_def = reinterpret_cast(src); + dst_bytes = dst_def; + src_bytes = src_def; + for (int j{8}; j > 1; j >>= 1) { cpy_size = size / j; - for (int i{thread_id}; i < cpy_size; i += WF_SIZE) { - store_asm(src_bytes, dst_bytes, j); + for (int i{wave_tid}; i < cpy_size; i += wave_size) { + dst_bytes = dst_def; + src_bytes = src_def; + src_bytes += i * j; dst_bytes += i * j; - size -= cpy_size * j; + + store_asm(src_bytes, dst_bytes, j); } + size -= cpy_size * j; + dst_def += cpy_size * j; + src_def += cpy_size * j; } if (size == 1) { diff --git a/tests/unit_tests/CMakeLists.txt b/tests/unit_tests/CMakeLists.txt index 7130c2c516..349bb2c2bf 100644 --- a/tests/unit_tests/CMakeLists.txt +++ b/tests/unit_tests/CMakeLists.txt @@ -71,6 +71,7 @@ target_sources( PRIVATE shmem_gtest.cpp heap_memory_gtest.cpp + hipmalloc_gtest.cpp bin_gtest.cpp binner_gtest.cpp #bitwise_gtest.cpp # Test is disabled becasue of compilation errors @@ -87,6 +88,8 @@ target_sources( notifier_gtest.cpp #forward_list_gtest.cpp free_list_gtest.cpp + #context_ipc_gtest.cpp + ipc_impl_simple_coarse_gtest.cpp ) ############################################################################### diff --git a/tests/unit_tests/context_ipc_gtest.cpp b/tests/unit_tests/context_ipc_gtest.cpp new file mode 100644 index 0000000000..c963a65673 --- /dev/null +++ b/tests/unit_tests/context_ipc_gtest.cpp @@ -0,0 +1,31 @@ +/****************************************************************************** + * 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 "context_ipc_gtest.hpp" + +using namespace rocshmem; + +TEST_F(ContextIpcTestFixture, constructor) { + /* do nothing for the moment, I *think* the + ** constructor is invoked automatically + */ +} diff --git a/tests/unit_tests/context_ipc_gtest.hpp b/tests/unit_tests/context_ipc_gtest.hpp new file mode 100644 index 0000000000..24a80fcedb --- /dev/null +++ b/tests/unit_tests/context_ipc_gtest.hpp @@ -0,0 +1,46 @@ +/****************************************************************************** + * 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 ROCSHMEM_CONTEXT_IPC_GTEST_HPP +#define ROCSHMEM_CONTEXT_IPC_GTEST_HPP + +#include "gtest/gtest.h" + +#include "../src/ipc/context_ipc_device.hpp" +#include "../src/ipc/backend_ipc.hpp" + +namespace rocshmem { + +class ContextIpcTestFixture : public ::testing::Test +{ + protected: + /** + * @brief Context Ipc Test + */ + IPCBackend be{MPI_COMM_WORLD}; + + IPCContext ipc_context_ {&be}; +}; + +} // namespace rocshmem + +#endif // ROCSHMEM_CONTEXT_IPC_GTEST_HPP diff --git a/tests/unit_tests/hipmalloc_gtest.cpp b/tests/unit_tests/hipmalloc_gtest.cpp new file mode 100644 index 0000000000..598952c0e6 --- /dev/null +++ b/tests/unit_tests/hipmalloc_gtest.cpp @@ -0,0 +1,43 @@ +/****************************************************************************** + * 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 "hipmalloc_gtest.hpp" + +using namespace rocshmem; + +TEST_F(HipMallocTestFixture, normal_1GBx256) { + void* ptr{nullptr}; + size_t gb {1073741824}; + for (int i{0}; i < 256; i++) { + hip_allocator_.allocate(&ptr, gb); + hip_allocator_.deallocate(ptr); + } +} + +TEST_F(HipMallocTestFixture, fine_1GBx256) { + void* ptr{nullptr}; + size_t gb {1073741824}; + for (int i{0}; i < 256; i++) { + hip_allocator_fg_.allocate(&ptr, gb); + hip_allocator_fg_.deallocate(ptr); + } +} diff --git a/tests/unit_tests/hipmalloc_gtest.hpp b/tests/unit_tests/hipmalloc_gtest.hpp new file mode 100644 index 0000000000..d6d341a57b --- /dev/null +++ b/tests/unit_tests/hipmalloc_gtest.hpp @@ -0,0 +1,41 @@ +/****************************************************************************** + * 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 ROCSHMEM_HIPMALLOC_GTEST_HPP +#define ROCSHMEM_HIPMALLOC_GTEST_HPP + +#include "gtest/gtest.h" + +#include "../src/memory/symmetric_heap.hpp" +#include "../src/util.hpp" + +namespace rocshmem { + +class HipMallocTestFixture : public ::testing::Test { + public: + HIPAllocator hip_allocator_ {}; + HIPAllocatorFinegrained hip_allocator_fg_ {}; +}; + +} // namespace rocshmem + +#endif // ROCSHMEM_HIPMALLOC_GTEST_HPP diff --git a/tests/unit_tests/ipc_impl_simple_coarse_gtest.cpp b/tests/unit_tests/ipc_impl_simple_coarse_gtest.cpp new file mode 100644 index 0000000000..dbf72923b5 --- /dev/null +++ b/tests/unit_tests/ipc_impl_simple_coarse_gtest.cpp @@ -0,0 +1,1029 @@ +/****************************************************************************** + * 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_impl_simple_coarse_gtest.hpp" + +using namespace rocshmem; + +TEST_F(IPCImplSimpleCoarseTestFixture, ptr_check) { + ASSERT_NE(heap_mem_.get_ptr(), nullptr); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, MPI_num_pes) { + ASSERT_EQ(mpi_.num_pes(), 2); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, IPC_bases) { + for(int i{0}; i < mpi_.num_pes(); i++) { + ASSERT_NE(ipc_impl_.ipc_bases[i], nullptr); + } +} + +TEST_F(IPCImplSimpleCoarseTestFixture, golden_1048576_int) { + iota_golden(1048576); + validate_golden(1048576); +} + +//============================================================================= + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wg_1x1x1_1024x1x1_32_int) { + dim3 grid {1,1,1}; + dim3 block {1024,1,1}; + write_wg(grid, block, 32); +} + +//============================================================================= + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wg_1x1x1_1024x1x1_32_int) { + dim3 grid {1,1,1}; + dim3 block {1024,1,1}; + read_wg(grid, block, 32); +} + +//============================================================================= + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wg_1x1x1_1x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {1,1,1}; + write_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wg_1x1x1_2x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {2,1,1}; + write_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wg_1x1x1_4x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {4,1,1}; + write_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wg_1x1x1_8x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {8,1,1}; + write_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wg_1x1x1_16x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {16,1,1}; + write_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wg_1x1x1_32x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {32,1,1}; + write_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wg_1x1x1_64x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {64,1,1}; + write_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wg_1x1x1_128x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {128,1,1}; + write_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wg_1x1x1_256x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {256,1,1}; + write_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wg_1x1x1_512x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {512,1,1}; + write_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wg_1x1x1_768x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {768,1,1}; + write_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wg_1x1x1_1024x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {1024,1,1}; + write_wg(grid, block, 1048576); +} + +//============================================================================= + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wg_1x1x1_1x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {1,1,1}; + read_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wg_1x1x1_2x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {2,1,1}; + read_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wg_1x1x1_4x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {4,1,1}; + read_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wg_1x1x1_8x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {8,1,1}; + read_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wg_1x1x1_16x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {16,1,1}; + read_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wg_1x1x1_32x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {32,1,1}; + read_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wg_1x1x1_64x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {64,1,1}; + read_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wg_1x1x1_128x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {128,1,1}; + read_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wg_1x1x1_256x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {256,1,1}; + read_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wg_1x1x1_512x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {512,1,1}; + read_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wg_1x1x1_768x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {768,1,1}; + read_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wg_1x1x1_1024x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {1024,1,1}; + read_wg(grid, block, 1048576); +} + +//============================================================================= + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_64x1x1_1_int) { + dim3 grid {1,1,1}; + dim3 block {64,1,1}; + write_wave(grid, block, 1); +} + +//============================================================================= + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_64x1x1_1_int) { + dim3 grid {1,1,1}; + dim3 block {64,1,1}; + read_wave(grid, block, 1); +} + +//============================================================================= + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_64x1x1_32_int) { + dim3 grid {1,1,1}; + dim3 block {64,1,1}; + write_wave(grid, block, 32); +} + +//============================================================================= + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_64x1x1_32_int) { + dim3 grid {1,1,1}; + dim3 block {64,1,1}; + read_wave(grid, block, 32); +} + +//============================================================================= + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_1x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {1,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_2x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {2,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_3x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {3,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_4x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {4,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_5x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {5,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_6x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {6,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_7x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {7,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_8x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {8,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_9x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {9,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_10x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {10,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_11x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {11,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_12x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {12,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_13x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {13,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_14x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {14,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_15x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {15,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_16x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {16,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_17x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {17,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_18x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {18,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_19x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {19,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_20x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {20,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_21x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {21,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_22x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {22,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_23x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {23,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_24x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {24,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_25x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {25,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_26x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {26,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_27x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {27,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_28x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {28,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_29x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {29,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_30x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {30,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_31x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {31,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_32x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {32,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_33x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {33,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_34x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {34,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_35x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {35,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_36x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {36,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_37x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {37,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_38x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {38,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_39x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {39,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_40x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {40,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_41x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {41,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_42x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {42,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_43x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {43,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_44x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {44,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_45x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {45,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_46x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {46,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_47x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {47,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_48x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {48,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_49x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {49,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_50x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {50,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_51x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {51,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_52x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {52,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_53x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {53,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_54x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {54,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_55x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {55,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_56x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {56,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_57x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {57,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_58x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {58,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_59x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {59,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_60x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {60,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_61x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {61,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_62x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {62,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_63x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {63,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, write_wave_1x1x1_64x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {64,1,1}; + write_wave(grid, block, 1048576); +} + +//============================================================================= + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_1x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {1,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_2x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {2,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_3x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {3,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_4x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {4,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_5x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {5,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_6x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {6,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_7x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {7,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_8x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {8,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_9x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {9,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_10x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {10,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_11x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {11,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_12x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {12,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_13x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {13,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_14x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {14,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_15x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {15,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_16x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {16,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_17x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {17,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_18x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {18,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_19x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {19,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_20x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {20,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_21x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {21,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_22x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {22,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_23x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {23,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_24x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {24,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_25x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {25,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_26x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {26,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_27x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {27,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_28x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {28,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_29x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {29,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_30x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {30,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_31x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {31,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_32x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {32,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_33x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {33,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_34x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {34,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_35x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {35,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_36x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {36,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_37x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {37,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_38x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {38,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_39x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {39,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_40x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {40,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_41x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {41,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_42x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {42,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_43x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {43,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_44x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {44,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_45x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {45,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_46x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {46,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_47x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {47,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_48x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {48,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_49x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {49,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_50x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {50,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_51x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {51,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_52x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {52,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_53x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {53,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_54x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {54,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_55x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {55,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_56x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {56,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_57x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {57,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_58x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {58,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_59x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {59,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_60x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {60,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_61x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {61,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_62x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {62,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_63x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {63,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleCoarseTestFixture, read_wave_1x1x1_64x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {64,1,1}; + read_wave(grid, block, 1048576); +} + +//============================================================================= + +TEST_F(IPCImplSimpleCoarseTestFixture, write_1x1x1_1x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {1,1,1}; + write(grid, block, 1048576); +} + +//============================================================================= + +TEST_F(IPCImplSimpleCoarseTestFixture, read_1x1x1_1x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {1,1,1}; + read(grid, block, 1048576); +} + diff --git a/tests/unit_tests/ipc_impl_simple_coarse_gtest.hpp b/tests/unit_tests/ipc_impl_simple_coarse_gtest.hpp new file mode 100644 index 0000000000..02dfd8c55a --- /dev/null +++ b/tests/unit_tests/ipc_impl_simple_coarse_gtest.hpp @@ -0,0 +1,236 @@ +/****************************************************************************** + * 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 ROCSHMEM_IPC_IMPL_SIMPLE_COARSE_GTEST_HPP +#define ROCSHMEM_IPC_IMPL_SIMPLE_COARSE_GTEST_HPP + +#include "gtest/gtest.h" + +#include + +#include +#include "../src/memory/symmetric_heap.hpp" +#include "../src/ipc_policy.hpp" + +namespace rocshmem { + +__global__ +void +kernel_simple_coarse_copy(IpcImpl *ipc_impl, int *src, int *dest, size_t bytes) { + if (!threadIdx.x) { + ipc_impl->ipcCopy(dest, src, bytes); + ipc_impl->ipcFence(); + } + __syncthreads(); +} + +__global__ +void +kernel_simple_coarse_copy_wg(IpcImpl *ipc_impl, int *src, int *dest, size_t bytes) { + ipc_impl->ipcCopy_wg(dest, src, bytes); + ipc_impl->ipcFence(); + __syncthreads(); +} + +__global__ +void +kernel_simple_coarse_copy_wave(IpcImpl *ipc_impl, int *src, int *dest, size_t bytes) { + ipc_impl->ipcCopy_wave(dest, src, bytes); + ipc_impl->ipcFence(); + __syncthreads(); +} + +class IPCImplSimpleCoarseTestFixture : public ::testing::Test { + + using HEAP_T = HeapMemory; + + using MPI_T = RemoteHeapInfo; + + using FN_T = void (*)(IpcImpl*, int*, int*, size_t); + + public: + IPCImplSimpleCoarseTestFixture() { + ipc_impl_.ipcHostInit(mpi_.my_pe(), mpi_.get_heap_bases() , MPI_COMM_WORLD); + + assert(ipc_impl_dptr_ == nullptr); + hip_allocator_.allocate((void**)&ipc_impl_dptr_, sizeof(IpcImpl)); + + CHECK_HIP(hipMemcpy(ipc_impl_dptr_, &ipc_impl_, + sizeof(IpcImpl), hipMemcpyHostToDevice)); + } + + ~IPCImplSimpleCoarseTestFixture() { + if (ipc_impl_dptr_) { + hip_allocator_.deallocate(ipc_impl_dptr_); + } + + ipc_impl_.ipcHostStop(); + } + + void launch(FN_T f, const dim3 grid, const dim3 block, int* src, int* dest, size_t bytes) { + f<<>>(ipc_impl_dptr_, src, dest, bytes); + CHECK_HIP(hipStreamSynchronize(nullptr)); + } + + enum TestType { + READ = 0, + WRITE = 1 + }; + + void write(const dim3 grid, const dim3 block, size_t elems) { + iota_golden(elems); + initialize_src_buffer(WRITE); + copy(WRITE, grid, block); + validate_dest_buffer(WRITE); + } + + void write_wg(const dim3 grid, const dim3 block, size_t elems) { + iota_golden(elems); + initialize_src_buffer(WRITE); + copy_wg(WRITE, grid, block); + validate_dest_buffer(WRITE); + } + + void write_wave(const dim3 grid, const dim3 block, size_t elems) { + iota_golden(elems); + initialize_src_buffer(WRITE); + copy_wave(WRITE, grid, block); + validate_dest_buffer(WRITE); + } + + void read(const dim3 grid, const dim3 block, size_t elems) { + iota_golden(elems); + initialize_src_buffer(READ); + copy(READ, grid, block); + validate_dest_buffer(READ); + } + + void read_wg(const dim3 grid, const dim3 block, size_t elems) { + iota_golden(elems); + initialize_src_buffer(READ); + copy_wg(READ, grid, block); + validate_dest_buffer(READ); + } + + void read_wave(const dim3 grid, const dim3 block, size_t elems) { + iota_golden(elems); + initialize_src_buffer(READ); + copy_wave(READ, grid, block); + validate_dest_buffer(READ); + } + + void iota_golden(size_t elems) { + golden_.resize(elems); + std::iota(golden_.begin(), golden_.end(), 0); + } + + void validate_golden(size_t elems) { + ASSERT_EQ(golden_.size(), elems); + for (int i{0}; i < golden_.size(); i++) { + ASSERT_EQ(golden_[i], i); + } + } + + void initialize_src_buffer(TestType test) { + if (!pe_initializes_src_buffer(test)) { + return; + } + size_t bytes = golden_.size() * sizeof(int); + auto dev_src = reinterpret_cast(ipc_impl_.ipc_bases[mpi_.my_pe()]); + CHECK_HIP(hipMemcpy(dev_src, golden_.data(), bytes, hipMemcpyHostToDevice)); + CHECK_HIP(hipStreamSynchronize(nullptr)); + } + + bool pe_initializes_src_buffer(TestType test) { + bool is_write_test = test; + bool is_read_test = !test; + return (is_write_test && mpi_.my_pe() == 0) || + (is_read_test && mpi_.my_pe() == 1); + } + + void execute(TestType test, FN_T fn, const dim3 grid, const dim3 block) { + if (mpi_.my_pe()) { + mpi_.barrier(); + mpi_.barrier(); + return; + } + int *src{nullptr}; + int *dest{nullptr}; + if (test == WRITE) { + src = reinterpret_cast(ipc_impl_.ipc_bases[0]); + dest = reinterpret_cast(ipc_impl_.ipc_bases[1]); + } else { + src = reinterpret_cast(ipc_impl_.ipc_bases[1]); + dest = reinterpret_cast(ipc_impl_.ipc_bases[0]); + } + size_t bytes = golden_.size() * sizeof(int); + mpi_.barrier(); + launch(fn, grid, block, src, dest, bytes); + mpi_.barrier(); + } + + void copy(TestType test, dim3 grid, dim3 block) { + execute(test, kernel_simple_coarse_copy, grid, block); + } + + void copy_wg(TestType test, dim3 grid, dim3 block) { + execute(test, kernel_simple_coarse_copy_wg, grid, block); + } + + void copy_wave(TestType test, dim3 grid, dim3 block) { + execute(test, kernel_simple_coarse_copy_wave, grid, block); + } + + void validate_dest_buffer(TestType test) { + if (!pe_validates_dest_buffer(test)) { + return; + } + + auto dev_dest = reinterpret_cast(ipc_impl_.ipc_bases[mpi_.my_pe()]); + for (int i{0}; i < golden_.size(); i++) { + ASSERT_EQ(golden_[i], dev_dest[i]); + } + } + + bool pe_validates_dest_buffer(TestType test) { + return !pe_initializes_src_buffer(test); + } + + protected: + std::vector golden_; + + std::vector output_; + + HEAP_T heap_mem_ {}; + + MPI_T mpi_ {heap_mem_.get_ptr(), heap_mem_.get_size()}; + + IpcImpl ipc_impl_ {}; + + IpcImpl *ipc_impl_dptr_ {nullptr}; + + HIPAllocator hip_allocator_ {}; +}; + +} // namespace rocshmem + +#endif // ROCSHMEM_IPC_IMPL_SIMPLE_COARSE_GTEST_HPP