diff --git a/projects/rocshmem/.gitignore b/projects/rocshmem/.gitignore new file mode 100644 index 0000000000..567609b123 --- /dev/null +++ b/projects/rocshmem/.gitignore @@ -0,0 +1 @@ +build/ diff --git a/projects/rocshmem/CMakeLists.txt b/projects/rocshmem/CMakeLists.txt index d068315c62..5cf74fb6cb 100644 --- a/projects/rocshmem/CMakeLists.txt +++ b/projects/rocshmem/CMakeLists.txt @@ -61,7 +61,6 @@ option(USE_IPC "Enable IPC support (using HIP)" OFF) option(USE_THREADS "Enable workgroup threads to share network queues" OFF) option(USE_WF_COAL "Enable wavefront message coalescing" OFF) option(USE_COHERENT_HEAP "Enable support for coherent systems" OFF) -option(USE_CACHED_HEAP "Enable support for cached systems" OFF) option(USE_MANAGED_HEAP "Enable managed memory" OFF) option(USE_HOST_HEAP "Enable host memory using malloc/free" OFF) option(USE_HIP_HOST_HEAP "Enable host memory using hip api" OFF) @@ -69,6 +68,7 @@ option(USE_FUNC_CALL "Force compiler to use function calls on library API" OFF) option(USE_SHARED_CTX "Request support for shared ctx between WG" OFF) option(USE_SINGLE_NODE "Enable single node support only." OFF) option(USE_HOST_SIDE_HDP_FLUSH "Use a polling thread to flush the HDP cache on the host." OFF) +option(USE_COOPERATIVE_GROUPS "Use cooperative groups for internal syncronization" OFF) option(BUILD_FUNCTIONAL_TESTS "Build the functional tests" ON) option(BUILD_SOS_TESTS "Build the host-facing tests" OFF) option(BUILD_UNIT_TESTS "Build the unit tests" ON) diff --git a/projects/rocshmem/README.md b/projects/rocshmem/README.md index c974fd7052..ba671ff51a 100644 --- a/projects/rocshmem/README.md +++ b/projects/rocshmem/README.md @@ -45,7 +45,7 @@ ROC_SHMEM base requirements: * May work with other versions, but not tested * AMD GFX9 GPUs (e.g.: MI25, Vega 56, Vega 64, MI50, MI60, MI100, Radeon VII) * AMD MI200 GPUs: To enable the support on MI200, please configure the library - with USE_CACHED_HEAP + with USE_COHERENT_HEAP * ROCm-aware MPI as described in [Building the Dependencies](#building-the-dependencies) * InfiniBand adaptor compatable with ROCm RDMA technology diff --git a/projects/rocshmem/cmake/config.h.in b/projects/rocshmem/cmake/config.h.in index 3b2732137a..21619a6c3e 100644 --- a/projects/rocshmem/cmake/config.h.in +++ b/projects/rocshmem/cmake/config.h.in @@ -8,10 +8,10 @@ #cmakedefine USE_SHARED_CTX #cmakedefine USE_WF_COAL #cmakedefine USE_COHERENT_HEAP -#cmakedefine USE_CACHED_HEAP #cmakedefine USE_MANAGED_HEAP #cmakedefine USE_HOST_HEAP #cmakedefine USE_HIP_HOST_HEAP #cmakedefine USE_FUNC_CALL #cmakedefine USE_SINGLE_NODE -#cmakedefine USE_HOST_SIDE_HDP_FLUSH \ No newline at end of file +#cmakedefine USE_HOST_SIDE_HDP_FLUSH +#cmakedefine USE_COOPERATIVE_GROUPS diff --git a/projects/rocshmem/scripts/build_configs/ipc_single_cg b/projects/rocshmem/scripts/build_configs/ipc_single_cg new file mode 100755 index 0000000000..c0904bad93 --- /dev/null +++ b/projects/rocshmem/scripts/build_configs/ipc_single_cg @@ -0,0 +1,31 @@ +#!/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 \ + -DUSE_COOPERATIVE_GROUPS=ON \ + $src_path +cmake --build . --parallel 8 +cmake --install . diff --git a/projects/rocshmem/scripts/build_configs/rc_single b/projects/rocshmem/scripts/build_configs/rc_single index 31eb7283c6..aceca7dbed 100755 --- a/projects/rocshmem/scripts/build_configs/rc_single +++ b/projects/rocshmem/scripts/build_configs/rc_single @@ -20,7 +20,6 @@ cmake \ -DUSE_DC=OFF \ -DUSE_IPC=OFF \ -DUSE_COHERENT_HEAP=OFF \ - -DUSE_CACHED_HEAP=OFF \ -DUSE_THREADS=OFF \ -DUSE_WF_COAL=OFF \ $src_path diff --git a/projects/rocshmem/scripts/build_configs/rc_single_single_node b/projects/rocshmem/scripts/build_configs/rc_single_single_node index 8bcf9fcb0d..a5223c379f 100755 --- a/projects/rocshmem/scripts/build_configs/rc_single_single_node +++ b/projects/rocshmem/scripts/build_configs/rc_single_single_node @@ -20,12 +20,10 @@ cmake \ -DUSE_DC=OFF \ -DUSE_IPC=ON \ -DUSE_COHERENT_HEAP=OFF \ - -DUSE_CACHED_HEAP=OFF \ -DUSE_THREADS=OFF \ -DUSE_WF_COAL=OFF \ -DUSE_SINGLE_NODE=ON \ -DUSE_HOST_SIDE_HDP_FLUSH=ON\ - -DROCM_PATH="/opt/rocm-5.4.2/"\ $src_path cmake --build . --parallel 8 cmake --install . diff --git a/projects/rocshmem/scripts/build_configs/rc_single_single_node_debug b/projects/rocshmem/scripts/build_configs/rc_single_single_node_debug index b535f1b853..c4e2477da9 100755 --- a/projects/rocshmem/scripts/build_configs/rc_single_single_node_debug +++ b/projects/rocshmem/scripts/build_configs/rc_single_single_node_debug @@ -20,7 +20,6 @@ cmake \ -DUSE_DC=OFF \ -DUSE_IPC=ON \ -DUSE_COHERENT_HEAP=OFF \ - -DUSE_CACHED_HEAP=OFF \ -DUSE_THREADS=OFF \ -DUSE_WF_COAL=OFF \ -DUSE_SINGLE_NODE=ON \ diff --git a/projects/rocshmem/scripts/build_configs/ro_ipc b/projects/rocshmem/scripts/build_configs/ro_ipc new file mode 100755 index 0000000000..d57404b88c --- /dev/null +++ b/projects/rocshmem/scripts/build_configs/ro_ipc @@ -0,0 +1,29 @@ +#!/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_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 . diff --git a/projects/rocshmem/scripts/build_configs/ro_net b/projects/rocshmem/scripts/build_configs/ro_net index c1d09194f7..6abc528ca6 100755 --- a/projects/rocshmem/scripts/build_configs/ro_net +++ b/projects/rocshmem/scripts/build_configs/ro_net @@ -22,7 +22,7 @@ cmake \ -DUSE_IPC=OFF \ -DUSE_THREADS=ON \ -DUSE_WF_COAL=OFF \ - -DUSE_COHERENT_HEAP=ON \ + -DUSE_COHERENT_HEAP=OFF \ $src_path cmake --build . --parallel 8 cmake --install . diff --git a/projects/rocshmem/scripts/build_configs/ro_net_debug b/projects/rocshmem/scripts/build_configs/ro_net_debug index 67c3f2d0a5..c2b8afdc00 100755 --- a/projects/rocshmem/scripts/build_configs/ro_net_debug +++ b/projects/rocshmem/scripts/build_configs/ro_net_debug @@ -21,7 +21,7 @@ cmake \ -DUSE_IPC=OFF \ -DUSE_THREADS=ON \ -DUSE_WF_COAL=OFF \ - -DUSE_COHERENT_HEAP=ON \ + -DUSE_COHERENT_HEAP=OFF \ $src_path cmake --build . --parallel 8 cmake --install . diff --git a/projects/rocshmem/scripts/functional_tests/driver.sh b/projects/rocshmem/scripts/functional_tests/driver.sh index 5e6bc53ae3..f93045fae8 100755 --- a/projects/rocshmem/scripts/functional_tests/driver.sh +++ b/projects/rocshmem/scripts/functional_tests/driver.sh @@ -121,9 +121,9 @@ case $2 in echo "amoinc_n2_w1_z1" ROC_SHMEM_MAX_NUM_CONTEXTS=1 mpirun -np 2 $1 -w 1 -z 1 -a 11 > $3/amoinc_n2_w1_z1.log check amoinc_n2_w1_z1 - echo "pingpong_n2_w1" - ROC_SHMEM_MAX_NUM_CONTEXTS=1 mpirun -np 2 $1 -w 1 -a 14 > $3/pingpong_n2_w1.log - check pingpong_n2_w1 + # echo "pingpong_n2_w1" + # ROC_SHMEM_MAX_NUM_CONTEXTS=1 mpirun -np 2 $1 -w 1 -a 14 > $3/pingpong_n2_w1.log + # check pingpong_n2_w1 echo "amoset_n2_w1_z1" ROC_SHMEM_MAX_NUM_CONTEXTS=1 mpirun -np 2 $1 -w 1 -z 1 -a 44 > $3/amoset_n2_w1_z1.log check amoset_n2_w1_z1 @@ -163,9 +163,9 @@ case $2 in echo "amoinc_n2_w8_z1" ROC_SHMEM_MAX_NUM_CONTEXTS=8 mpirun -np 2 $1 -w 8 -z 1 -a 11 > $3/amoinc_n2_w8_z1.log check amoinc_n2_w8_z1 - echo "pingpong_n2_w1" - ROC_SHMEM_MAX_NUM_CONTEXTS=1 mpirun -np 2 $1 -w 1 -a 14 > $3/pingpong_n2_w1.log - check pingpong_n2_w1 + # echo "pingpong_n2_w1" + # ROC_SHMEM_MAX_NUM_CONTEXTS=1 mpirun -np 2 $1 -w 1 -a 14 > $3/pingpong_n2_w1.log + # check pingpong_n2_w1 echo "amoset_n2_w8_z1" ROC_SHMEM_MAX_NUM_CONTEXTS=8 mpirun -np 2 $1 -w 8 -z 1 -a 44 > $3/amoset_n2_w8_z1.log check amoset_n2_w8_z1 diff --git a/projects/rocshmem/scripts/unit_tests/driver.sh b/projects/rocshmem/scripts/unit_tests/driver.sh new file mode 100755 index 0000000000..8eea7daef7 --- /dev/null +++ b/projects/rocshmem/scripts/unit_tests/driver.sh @@ -0,0 +1,62 @@ +#!/bin/bash + +# Function to display help information +function display_help { + echo "Usage:" + echo " $0 binary_name all # Runs all standard tests" + echo " $0 binary_name custom # Runs custom test configuration" + echo + echo "Arguments:" + echo " binary_name: Name of the binary to run." + echo " all: Executes predefined test configurations." + echo " custom: Executes a test with custom MPI ranks and GTest filter." + echo " ranks: Number of MPI ranks (required for custom mode)." + echo " filter: GTest filter string (required for custom mode)." + echo +} + +# Validate number of arguments for each mode +if [[ "$#" -lt 2 ]] || + { [[ "$2" == "all" ]] && [[ "$#" -ne 2 ]]; } || + { [[ "$2" == "custom" ]] && [[ "$#" -ne 4 ]]; }; then + display_help + exit 1 +fi + +binary_name=$1 +mode=$2 +timestamp=$(date "+%Y-%m-%d-%H:%M:%S") +log_file="unit_tests_${timestamp}.log" + +# Function to execute mpirun command +function run_mpirun { + local np=$1 + local gtest_filter=$2 + echo "mpirun -np $np $binary_name --gtest_filter='$gtest_filter'" | tee -a "$log_file" + mpirun -np "$np" "$binary_name" --gtest_filter="$gtest_filter" >> "$log_file" 2>&1 +} + +# Processing modes +case $mode in + all) + run_mpirun 4 "-IPCImplSimpleCoarseTestFixture.*:IPCImplSimpleFineTestFixture.*" + run_mpirun 2 "IPCImplSimpleCoarseTestFixture.*:IPCImplSimpleFineTestFixture.*" + ;; + custom) + # Check if ranks is a positive integer + if [[ "$3" -le 1 ]]; then + echo "Error: 'ranks' must be a positive integer." + display_help + exit 1 + fi + run_mpirun $3 $4 + ;; + *) + echo "Error: Invalid mode '$mode'." | tee -a "$log_file" + display_help + exit 1 + ;; +esac + +echo "Tests Completed" +echo "log file: '$log_file'" diff --git a/projects/rocshmem/src/assembly.hpp b/projects/rocshmem/src/assembly.hpp new file mode 100644 index 0000000000..a14907a0bd --- /dev/null +++ b/projects/rocshmem/src/assembly.hpp @@ -0,0 +1,271 @@ +/****************************************************************************** + * 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_ASSEMBLY_HPP_ +#define LIBRARY_SRC_ASSEMBLY_HPP_ + +#include +#include +#include + +namespace rocshmem { + +#define DO_PRAGMA(x) _Pragma(#x) +#define NOWARN(warnoption, ...) \ + DO_PRAGMA(GCC diagnostic push) \ + DO_PRAGMA(GCC diagnostic ignored #warnoption) \ + __VA_ARGS__ \ + DO_PRAGMA(GCC diagnostic pop) + +#define SFENCE() asm volatile("sfence" ::: "memory") + +__device__ __forceinline__ int uncached_load_ubyte(uint8_t* src) { + int ret; +#if defined(__gfx906__) +#endif +#if defined(__gfx908__) +#endif +#if defined(__gfx90a__) + asm volatile( + "global_load_ubyte %0 %1 off glc slc \n" + "s_waitcnt vmcnt(0)" + : "=v"(ret) + : "v"(src)); +#endif +#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__) + asm volatile( + "global_load_ubyte %0 %1 off sc0 sc1 \n" + "s_waitcnt vmcnt(0)" + : "=v"(ret) + : "v"(src)); +#endif + return ret; +} + +__device__ __forceinline__ void refresh_volatile_sbyte(volatile int *assigned_value, + volatile char *read_value) { +#if defined(__gfx906__) +#endif +#if defined(__gfx908__) +#endif +#if defined(__gfx90a__) + asm volatile( + "global_load_sbyte %0 %1 off glc slc\n " + "s_waitcnt vmcnt(0)" + : "=v"(*assigned_value) + : "v"(read_value)); +#endif +#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__) + asm volatile( + "global_load_sbyte %0 %1 off sc0 sc1\n " + "s_waitcnt vmcnt(0)" + : "=v"(*assigned_value) + : "v"(read_value)); +#endif +} + +__device__ __forceinline__ void refresh_volatile_dwordx2(volatile uint64_t *assigned_value, + volatile uint64_t *read_value) { +#if defined(__gfx906__) +#endif +#if defined(__gfx908__) +#endif +#if defined(__gfx90a__) + asm volatile( + "global_load_dwordx2 %0 %1 off glc slc\n " + "s_waitcnt vmcnt(0)" + : "=v"(*assigned_value) + : "v"(read_value)); +#endif +#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__) + asm volatile( + "global_load_dwordx2 %0 %1 off sc0 sc1\n " + "s_waitcnt vmcnt(0)" + : "=v"(*assigned_value) + : "v"(read_value)); +#endif +} + +/* Ignore the warning about deprecated volatile. + * The only usage of volatile is to force the compiler to generate + * the assembly instruction. If volatile is omitted, the compiler + * will NOT generate the non-temporal load or the waitcnt. + */ +// clang-format off +NOWARN(-Wdeprecated-volatile, + template __device__ __forceinline__ T uncached_load(T* src) { + T ret; + switch (sizeof(T)) { + case 4: +#if defined(__gfx906__) +#endif +#if defined(__gfx908__) +#endif +#if defined(__gfx90a__) + asm volatile( + "global_load_dword %0 %1 off glc slc \n" + "s_waitcnt vmcnt(0)" + : "=v"(ret) + : "v"(src)); +#endif +#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__) + asm volatile( + "global_load_dword %0 %1 off sc0 sc1 \n" + "s_waitcnt vmcnt(0)" + : "=v"(ret) + : "v"(src)); +#endif + break; + case 8: +#if defined(__gfx906__) +#endif +#if defined(__gfx908__) +#endif +#if defined(__gfx90a__) + asm volatile( + "global_load_dwordx2 %0 %1 off glc slc \n" + "s_waitcnt vmcnt(0)" + : "=v"(ret) + : "v"(src)); +#endif +#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__) + asm volatile( + "global_load_dwordx2 %0 %1 off sc0 sc1 \n" + "s_waitcnt vmcnt(0)" + : "=v"(ret) + : "v"(src)); +#endif + break; + default: + break; + } + return ret; + } +) +// clang-format on + +__device__ __forceinline__ void __roc_inv() { +#if defined USE_COHERENT_HEAP +#if defined(__gfx906__) +#endif +#if defined(__gfx908__) +#endif +#if defined(__gfx90a__) +// asm volatile("buffer_wbinvl1;"); +#endif +#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__) +// asm volatile("buffer_inv sc0 sc1;"); +#endif +#endif +} + +__device__ __forceinline__ void __roc_flush() { +#if defined USE_COHERENT_HEAP +#if defined(__gfx906__) +#endif +#if defined(__gfx908__) +#endif +#if defined(__gfx90a__) +// asm volatile("s_dcache_wb;"); +// asm volatile("buffer_wbl2;"); +#endif +#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__) +// asm volatile("s_dcache_wb;"); +// asm volatile("buffer_wbl2;"); +#endif +#endif +} + +__device__ __forceinline__ void store_asm(uint8_t* val, uint8_t* dst, + int size) { + switch (size) { + case 2: { + int16_t val16{*(reinterpret_cast(val))}; +#if defined(__gfx906__) +#endif +#if defined(__gfx908__) +#endif +#if defined(__gfx90a__) + asm volatile("flat_store_short %0 %1 glc slc" : : "v"(dst), "v"(val16)); +#endif +#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__) + asm volatile("flat_store_short %0 %1 sc0 sc1" : : "v"(dst), "v"(val16)); +#endif + break; + } + case 4: { + int32_t val32{*(reinterpret_cast(val))}; +#if defined(__gfx906__) +#endif +#if defined(__gfx908__) +#endif +#if defined(__gfx90a__) + asm volatile("flat_store_dword %0 %1 glc slc" : : "v"(dst), "v"(val32)); +#endif +#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__) + asm volatile("flat_store_dword %0 %1 sc0 sc1" : : "v"(dst), "v"(val32)); +#endif + break; + } + case 8: { + int64_t val64{*(reinterpret_cast(val))}; +#if defined(__gfx906__) +#endif +#if defined(__gfx908__) +#endif +#if defined(__gfx90a__) + asm volatile("flat_store_dwordx2 %0 %1 glc slc" : : "v"(dst), "v"(val64)); +#endif +#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__) + asm volatile("flat_store_dwordx2 %0 %1 sc0 sc1" : : "v"(dst), "v"(val64)); +#endif + break; + } + default: + break; + } +} + +__device__ __forceinline__ uint64_t __read_clock() { + uint64_t clock{}; +#if defined(__gfx906__) +#endif +#if defined(__gfx908__) +#endif +#if defined(__gfx90a__) + asm volatile( + "s_memrealtime %0\n" + "s_waitcnt lgkmcnt(0)\n" + : "=s"(clock)); +#endif +#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__) + asm volatile( + "s_memrealtime %0\n" + "s_waitcnt lgkmcnt(0)\n" + : "=s"(clock)); +#endif + return clock; +} + +} // namespace rocshmem + +#endif // LIBRARY_SRC_ASSEMBLY_HPP_ diff --git a/projects/rocshmem/src/atomic.hpp b/projects/rocshmem/src/atomic.hpp new file mode 100644 index 0000000000..f0828e4145 --- /dev/null +++ b/projects/rocshmem/src/atomic.hpp @@ -0,0 +1,141 @@ +/****************************************************************************** + * 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_ATOMIC_HPP +#define LIBRARY_SRC_ATOMIC_HPP + +#include + +namespace rocshmem { +namespace detail { +namespace atomic { + +typedef enum rocshmem_memory_scope { + memory_scope_thread = __HIP_MEMORY_SCOPE_SINGLETHREAD, + memory_scope_wavefront = __HIP_MEMORY_SCOPE_WAVEFRONT, + memory_scope_workgroup = __HIP_MEMORY_SCOPE_WORKGROUP, + memory_scope_agent = __HIP_MEMORY_SCOPE_AGENT, + memory_scope_system = __HIP_MEMORY_SCOPE_SYSTEM, +} rocshmem_memory_scope; + +typedef enum rocshmem_memory_order { + memory_order_relaxed = __ATOMIC_RELAXED, + memory_order_consume = __ATOMIC_CONSUME, + memory_order_acquire = __ATOMIC_ACQUIRE, + memory_order_release = __ATOMIC_RELEASE, + memory_order_acq_rel = __ATOMIC_ACQ_REL, + memory_order_seq_cst = __ATOMIC_SEQ_CST +} rocshmem_memory_order; + +struct rocshmem_memory_orders { + rocshmem_memory_order load {memory_order_acquire}; + rocshmem_memory_order store {memory_order_release}; + rocshmem_memory_order atomic {memory_order_acq_rel}; + rocshmem_memory_order weak_cas_success {memory_order_acq_rel}; + rocshmem_memory_order weak_cas_failure {memory_order_acq_rel}; + rocshmem_memory_order strong_cas_success {memory_order_acq_rel}; + rocshmem_memory_order strong_cas_failure {memory_order_acq_rel}; +}; + +template +__host__ __device__ +T load(const T* address, rocshmem_memory_orders o) { + return __hip_atomic_load(address, o.load, s); +} + +template +__host__ __device__ +void store(T* address, const T value, rocshmem_memory_orders o) { + return __hip_atomic_store(address, value, o.store, s); +} + +template +__host__ __device__ +bool compare_exchange_weak(T& expected, T desired, rocshmem_memory_orders o) { + return __hip_atomic_compare_exchange_weak(expected, desired, o.weak_cas_success, o.weak_cas_failure, s); +} + +template +__host__ __device__ +bool compare_exchange_strong(T& expected, T desired, rocshmem_memory_orders o) { + return __hip_atomic_compare_exchange_strong(expected, desired, o.strong_cas_success, o.strong_cas_failure, s); +} + +template +__host__ __device__ +T fetch_add(T* obj, U arg, rocshmem_memory_orders o) { + return __hip_atomic_fetch_add(obj, arg, o.atomic, s); +} + +template +__host__ __device__ +T fetch_sub(T* obj, U arg, rocshmem_memory_orders o) { + return __hip_atomic_fetch_sub(obj, arg, o.atomic, s); +} + +template +__host__ __device__ +T fetch_and(T* obj, U arg, rocshmem_memory_orders o) { + return __hip_atomic_fetch_and(obj, arg, o.atomic, s); +} + +template +__host__ __device__ +T fetch_or(T* obj, U arg, rocshmem_memory_orders o) { + return __hip_atomic_fetch_or(obj, arg, o, s); +} + +template +__host__ __device__ +T fetch_xor(T* obj, U arg, rocshmem_memory_orders o) { + return __hip_atomic_fetch_xor(obj, arg, o.atomic, s); +} + +template +__host__ __device__ +T fetch_max(T* obj, U arg, rocshmem_memory_orders o) { + return __hip_atomic_fetch_max(obj, arg, o.atomic, s); +} + +template +__host__ __device__ +T fetch_min(T* obj, U arg, rocshmem_memory_orders o) { + return __hip_atomic_fetch_min(obj, arg, o.atomic, s); +} + +template +__device__ +void threadfence() { + if constexpr (s == memory_scope_system) { + __threadfence_system(); + } else if constexpr (s == memory_scope_agent) { + __threadfence(); + } else if constexpr (s == memory_scope_workgroup) { + __threadfence_block(); + } +} + +} // namespace atomic +} // namespace detail +} // namespace rocshmem + +#endif // LIBRARY_SRC_ATOMIC_HPP_ diff --git a/projects/rocshmem/src/backend_type.hpp b/projects/rocshmem/src/backend_type.hpp index d8b2972504..50818eb178 100644 --- a/projects/rocshmem/src/backend_type.hpp +++ b/projects/rocshmem/src/backend_type.hpp @@ -70,13 +70,11 @@ enum class BackendType { RO_BACKEND, GPU_IB_BACKEND, IPC_BACKEND }; */ #ifdef USE_GPU_IB #define DISPATCH_RET(Func) \ - auto ret_val{0}; \ - ret_val = static_cast(this)->Func; \ + auto ret_val = static_cast(this)->Func; \ return ret_val; #elif defined(USE_RO) #define DISPATCH_RET(Func) \ - auto ret_val{0}; \ - ret_val = static_cast(this)->Func; \ + auto ret_val = static_cast(this)->Func; \ return ret_val; #else #define DISPATCH_RET(Func) \ @@ -128,13 +126,11 @@ enum class BackendType { RO_BACKEND, GPU_IB_BACKEND, IPC_BACKEND }; #ifdef USE_GPU_IB #define HOST_DISPATCH_RET(Func) \ - auto ret_val{0}; \ - ret_val = static_cast(this)->Func; \ + auto ret_val = static_cast(this)->Func; \ return ret_val; #elif defined(USE_RO) #define HOST_DISPATCH_RET(Func) \ - auto ret_val{0}; \ - ret_val = static_cast(this)->Func; \ + auto ret_val = static_cast(this)->Func; \ return ret_val; #else #define HOST_DISPATCH_RET(Func) \ diff --git a/projects/rocshmem/src/gpu_ib/backend_ib.cpp b/projects/rocshmem/src/gpu_ib/backend_ib.cpp index 9ae3acd482..4692235738 100644 --- a/projects/rocshmem/src/gpu_ib/backend_ib.cpp +++ b/projects/rocshmem/src/gpu_ib/backend_ib.cpp @@ -133,8 +133,6 @@ __device__ bool GPUIBBackend::create_ctx(int64_t options, return false; } ctx_ = pop_result.value; - - ctx_->dev_mtx_.shareable_ = static_cast(options >> 3); ctx->ctx_opaque = ctx_; return true; } @@ -296,7 +294,8 @@ void GPUIBBackend::setup_default_ctx() { CHECK_HIP(hipGetSymbolAddress(reinterpret_cast(&symbol_address), HIP_SYMBOL(ROC_SHMEM_CTX_DEFAULT))); - roc_shmem_ctx_t ctx_default_host{default_ctx_, nullptr}; + TeamInfo *tinfo = team_tracker.get_team_world()->tinfo_wrt_world; + roc_shmem_ctx_t ctx_default_host{default_ctx_, tinfo}; hipStream_t stream; CHECK_HIP(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking)); diff --git a/projects/rocshmem/src/gpu_ib/connection.cpp b/projects/rocshmem/src/gpu_ib/connection.cpp index 631f72f324..b6b2ae4dcd 100644 --- a/projects/rocshmem/src/gpu_ib/connection.cpp +++ b/projects/rocshmem/src/gpu_ib/connection.cpp @@ -256,7 +256,7 @@ void* Connection::buf_alloc([[maybe_unused]] struct ibv_pd* pd, if (use_gpu_mem) { void* dev_ptr; if (coherent_cq == 1) { -#if defined USE_COHERENT_HEAP || defined USE_CACHED_HEAP +#if defined USE_COHERENT_HEAP CHECK_HIP(hipMalloc(reinterpret_cast(&dev_ptr), size)); #else #ifdef HIP_SUPPORTS_MALLOC_UNCACHED diff --git a/projects/rocshmem/src/gpu_ib/segment_builder.cpp b/projects/rocshmem/src/gpu_ib/segment_builder.cpp index 205c890861..6f2a77a23e 100644 --- a/projects/rocshmem/src/gpu_ib/segment_builder.cpp +++ b/projects/rocshmem/src/gpu_ib/segment_builder.cpp @@ -22,6 +22,7 @@ #include "segment_builder.hpp" +#include "../util.hpp" #include "endian.hpp" namespace rocshmem { @@ -35,9 +36,9 @@ __device__ SegmentBuilder::SegmentBuilder(uint64_t wqe_idx, void *base) { __device__ void SegmentBuilder::update_cntrl_seg( uint8_t opcode, uint16_t wqe_idx, uint32_t ctrl_qp_sq, uint64_t ctrl_sig, ConnectionImpl *connection_policy, bool zero_byte_rd) { - mlx5_wqe_ctrl_seg *ctrl_seg = &seg_ptr->ctrl_seg; + mlx5_wqe_ctrl_seg ctrl_seg; - ctrl_seg->opmod_idx_opcode = (opcode << 24) | (wqe_idx << 8); + ctrl_seg.opmod_idx_opcode = (opcode << 24) | (wqe_idx << 8); uint32_t DS = 2; if (zero_byte_rd == false) { @@ -48,39 +49,43 @@ __device__ void SegmentBuilder::update_cntrl_seg( DS += connection_policy->wqeCntrlOffset(); - ctrl_seg->qpn_ds = (DS << 24) | ctrl_qp_sq; + ctrl_seg.qpn_ds = (DS << 24) | ctrl_qp_sq; - ctrl_seg->signature = ctrl_sig; + ctrl_seg.signature = ctrl_sig; - ctrl_seg->fm_ce_se = ctrl_sig >> 24; + ctrl_seg.fm_ce_se = ctrl_sig >> 24; - ctrl_seg->imm = ctrl_sig >> 32; + ctrl_seg.imm = ctrl_sig >> 32; + + memcpy(&seg_ptr->ctrl_seg, &ctrl_seg, sizeof(mlx5_wqe_ctrl_seg)); seg_ptr++; } __device__ void SegmentBuilder::update_atomic_data_seg(uint64_t atomic_data, uint64_t atomic_cmp) { - mlx5_wqe_atomic_seg *atomic_seg = &seg_ptr->atomic_seg; + mlx5_wqe_atomic_seg atomic_seg; - swap_endian_store(reinterpret_cast(&atomic_seg->swap_add), + swap_endian_store(reinterpret_cast(&atomic_seg.swap_add), atomic_data); - swap_endian_store(reinterpret_cast(&atomic_seg->compare), + swap_endian_store(reinterpret_cast(&atomic_seg.compare), atomic_cmp); + memcpy(&seg_ptr->atomic_seg, &atomic_seg, sizeof(mlx5_wqe_atomic_seg)); seg_ptr++; } __device__ void SegmentBuilder::update_rdma_seg(uintptr_t *raddr, uint32_t rkey) { - mlx5_wqe_raddr_seg *raddr_seg = &seg_ptr->raddr_seg; + mlx5_wqe_raddr_seg raddr_seg; - raddr_seg->rkey = rkey; + raddr_seg.rkey = rkey; - swap_endian_store(reinterpret_cast(&raddr_seg->raddr), + swap_endian_store(reinterpret_cast(&raddr_seg.raddr), reinterpret_cast(raddr)); + memcpy(&seg_ptr->raddr_seg, &raddr_seg, sizeof(mlx5_wqe_raddr_seg)); seg_ptr++; } @@ -90,33 +95,36 @@ __device__ void SegmentBuilder::update_data_seg(uintptr_t *laddr, int32_t size, return; } - mlx5_wqe_data_seg *data_seg = &seg_ptr->data_seg; + mlx5_wqe_data_seg data_seg; + data_seg.lkey = lkey; - data_seg->lkey = lkey; - - swap_endian_store(&data_seg->byte_count, size & 0x7FFFFFFFU); - - swap_endian_store(reinterpret_cast(&data_seg->addr), + swap_endian_store(&data_seg.byte_count, size & 0x7FFFFFFFU); + swap_endian_store(reinterpret_cast(&data_seg.addr), reinterpret_cast(laddr)); + memcpy(&seg_ptr->data_seg, &data_seg, sizeof(mlx5_wqe_data_seg)); seg_ptr++; } __device__ void SegmentBuilder::update_inl_data_seg(uintptr_t *laddr, int32_t size) { - mlx5_wqe_inl_data_seg *inl_data_seg = &seg_ptr->inl_data_seg; + mlx5_wqe_inl_data_seg inl_data_seg; - swap_endian_store(&inl_data_seg->byte_count, (size & 0x3FF) | 0x80000000); + swap_endian_store(&inl_data_seg.byte_count, (size & 0x3FF) | 0x80000000); // Assume fence HDP flush // TODO(khamidou): Rework fence interface to avoid this + size_t field_size{sizeof(mlx5_wqe_inl_data_seg)}; if (!laddr) { uint8_t flush_val = 1; - memcpy(inl_data_seg + 1, &flush_val, sizeof(flush_val)); + memcpy(&inl_data_seg + 1, &flush_val, sizeof(flush_val)); + field_size += sizeof(flush_val); } else { - memcpy(inl_data_seg + 1, laddr, size); + memcpy(&inl_data_seg + 1, laddr, size); + field_size += size; } + memcpy(&seg_ptr->inl_data_seg, &inl_data_seg, field_size); seg_ptr++; } diff --git a/projects/rocshmem/src/ipc/CMakeLists.txt b/projects/rocshmem/src/ipc/CMakeLists.txt index e8d64ff523..07a40538a5 100644 --- a/projects/rocshmem/src/ipc/CMakeLists.txt +++ b/projects/rocshmem/src/ipc/CMakeLists.txt @@ -30,4 +30,5 @@ target_sources( context_ipc_host.cpp backend_ipc.cpp ipc_team.cpp + context_ipc_device_coll.cpp ) diff --git a/projects/rocshmem/src/ipc/backend_ipc.cpp b/projects/rocshmem/src/ipc/backend_ipc.cpp index 58caaf53f7..7475fb2c8b 100644 --- a/projects/rocshmem/src/ipc/backend_ipc.cpp +++ b/projects/rocshmem/src/ipc/backend_ipc.cpp @@ -56,18 +56,18 @@ int get_ls_non_zero_bit(char *bitmask, int mask_length) { 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 */ @@ -82,16 +82,19 @@ IPCBackend::IPCBackend(MPI_Comm comm) allocate_atomic_region(&bp->atomic_ret, MAX_NUM_BLOCKS); - default_context_proxy_ = IPCDefaultContextProxyT(this); - setup_team_world(); + TeamInfo *tinfo = team_tracker.get_team_world()->tinfo_wrt_world; + roc_shmem_collective_init(); + setup_fence_buffer(); + teams_init(); + default_context_proxy_ = IPCDefaultContextProxyT(this, tinfo); + setup_ctxs(); - } IPCBackend::~IPCBackend() { @@ -105,7 +108,7 @@ IPCBackend::~IPCBackend() { * Free the atomic_ret array. */ CHECK_HIP(hipFree(bp->atomic_ret->atomic_base_ptr)); - + // TODO(Avinash) Free g_ret // delete host_interface; @@ -141,6 +144,8 @@ __device__ bool IPCBackend::create_ctx(int64_t options, roc_shmem_ctx_t *ctx) { ctx_ = pop_result.value; ctx->ctx_opaque = ctx_; + + ctx_->tinfo = reinterpret_cast(ctx->team_opaque); return true; } @@ -287,6 +292,14 @@ void IPCBackend::teams_destroy() { free(reduced_bitmask_); } +void IPCBackend::setup_fence_buffer() { + /* + * Allocate heap space for fence + */ + fence_pool = reinterpret_cast(roc_shmem_malloc( + sizeof(int) * num_pes)); +} + void IPCBackend::roc_shmem_collective_init() { /* * Allocate heap space for barrier_sync @@ -389,4 +402,4 @@ void IPCBackend::teams_init() { NET_CHECK(MPI_Barrier(thread_comm)); } -} // namespace rocshmem \ No newline at end of file +} // namespace rocshmem diff --git a/projects/rocshmem/src/ipc/backend_ipc.hpp b/projects/rocshmem/src/ipc/backend_ipc.hpp index cc385f16ee..de8cb9c617 100644 --- a/projects/rocshmem/src/ipc/backend_ipc.hpp +++ b/projects/rocshmem/src/ipc/backend_ipc.hpp @@ -32,6 +32,10 @@ #include "ipc_context_proxy.hpp" #include "../ipc_policy.hpp" +#ifndef USE_COOPERATIVE_GROUPS +#include "../memory/notifier.hpp" +#endif /* NOT DEFINED: USE_COOPERATIVE_GROUPS */ + namespace rocshmem { class IPCBackend : public Backend { @@ -104,7 +108,7 @@ class IPCBackend : public Backend { 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) */ @@ -156,6 +160,15 @@ class IPCBackend : public Backend { */ void *pAta_pool{nullptr}; + /** + * @brief Handle for raw memory for fence/quiet + */ + int *fence_pool{nullptr}; + +#ifndef USE_COOPERATIVE_GROUPS + NotifierProxy notifier_{}; +#endif /* NOT DEFINED: USE_COOPERATIVE_GROUPS */ + protected: /** * @copydoc Backend::dump_backend_stats() @@ -203,6 +216,11 @@ class IPCBackend : public Backend { */ void roc_shmem_collective_init(); + /** + * @brief Allocate buffer for fence/quiet operation + */ + void setup_fence_buffer(); + private: /** * @brief Proxy for the default context @@ -249,4 +267,4 @@ class IPCBackend : public Backend { } // namespace rocshmem -#endif // LIBRARY_SRC_IPC_BACKEND_HPP_ \ No newline at end of file +#endif // LIBRARY_SRC_IPC_BACKEND_HPP_ diff --git a/projects/rocshmem/src/ipc/context_ipc_device.cpp b/projects/rocshmem/src/ipc/context_ipc_device.cpp index 4bf7072aa5..d208a070fb 100644 --- a/projects/rocshmem/src/ipc/context_ipc_device.cpp +++ b/projects/rocshmem/src/ipc/context_ipc_device.cpp @@ -44,8 +44,16 @@ __host__ IPCContext::IPCContext(Backend *b) 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; + fence_pool = backend->fence_pool; + + orders_.store = detail::atomic::rocshmem_memory_order::memory_order_seq_cst; + +#ifndef USE_COOPERATIVE_GROUPS + notifier_ = backend->notifier_.get(); +#endif /* NOT DEFINED: USE_COOPERATIVE_GROUPS */ } __device__ void IPCContext::threadfence_system() { @@ -84,12 +92,17 @@ __device__ void IPCContext::getmem_nbi(void *dest, const void *source, } __device__ void IPCContext::fence() { + for (int i{0}, j{tinfo->pe_start}; i < tinfo->size; i++, j += tinfo->stride) { + detail::atomic::store(&fence_pool[j], 1, orders_); + } } __device__ void IPCContext::fence(int pe) { + detail::atomic::store(&fence_pool[pe], 1, orders_); } __device__ void IPCContext::quiet() { + fence(); } __device__ void *IPCContext::shmem_ptr(const void *dest, int pe) { @@ -97,18 +110,6 @@ __device__ void *IPCContext::shmem_ptr(const void *dest, int pe) { return ret; } -__device__ void IPCContext::barrier_all() { - __syncthreads(); -} - -__device__ void IPCContext::sync_all() { - __syncthreads(); -} - -__device__ void IPCContext::sync(roc_shmem_team_t team) { - __syncthreads(); -} - __device__ void IPCContext::putmem_wg(void *dest, const void *source, size_t nelems, int pe) { uint64_t L_offset = diff --git a/projects/rocshmem/src/ipc/context_ipc_device.hpp b/projects/rocshmem/src/ipc/context_ipc_device.hpp index 377790ad67..065d37cec4 100644 --- a/projects/rocshmem/src/ipc/context_ipc_device.hpp +++ b/projects/rocshmem/src/ipc/context_ipc_device.hpp @@ -24,6 +24,15 @@ #define LIBRARY_SRC_IPC_CONTEXT_DEVICE_HPP_ #include "../context.hpp" +#include "../atomic.hpp" +#include "../team.hpp" + +#ifdef USE_COOPERATIVE_GROUPS +#include +namespace cg = cooperative_groups; +#else +#include "../memory/notifier.hpp" +#endif /* USE_COOPERATIVE_GROUPS */ namespace rocshmem { @@ -194,50 +203,62 @@ class IPCContext : public Context { template __device__ void get_nbi_wave(T *dest, const T *source, size_t nelems, int pe); - // Wait / Test functions - template - __device__ void wait_until(T* ptr, roc_shmem_cmps cmp, T val); - - template - __device__ void wait_until_all(T* ptr, size_t nelems, - const int *status, - roc_shmem_cmps cmp, T val); - - template - __device__ size_t wait_until_any(T* ptr, size_t nelems, - const int *status, - roc_shmem_cmps cmp, T val); - - template - __device__ size_t wait_until_some(T* ptr, size_t nelems, - size_t* indices, - const int *status, - roc_shmem_cmps cmp, T val); - - template - __device__ void wait_until_all_vector(T* ptr, size_t nelems, - const int *status, - roc_shmem_cmps cmp, T* vals); - - template - __device__ size_t wait_until_any_vector(T* ptr, size_t nelems, - const int *status, - roc_shmem_cmps cmp, T* vals); - template - __device__ size_t wait_until_some_vector(T* ptr, size_t nelems, - size_t* indices, - const int *status, - roc_shmem_cmps cmp, T* vals); - - template - __device__ int test(T* ptr, roc_shmem_cmps cmp, T val); - private: - uint64_t* atomic_base_ptr{nullptr}; + //context class has IpcImpl object (ipcImpl_) + IpcImpl *ipcImpl{nullptr}; - char* g_ret; + 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}; + +#ifndef USE_COOPERATIVE_GROUPS + Notifier *notifier_{nullptr}; +#endif /* NOT DEFINED: USE_COOPERATIVE_GROUPS */ + + //Struct defining memory ordering for atomic operations. + detail::atomic::rocshmem_memory_orders orders_{}; + + //Buffer to perform Atomic store to enforce memory ordering + int *fence_pool{nullptr}; + + public: + //TODO(Avinash): + //Make tinfo private variable, it requires changes to the context + //creation API in backend + + //Team information for the team associated with the context + TeamInfo *tinfo{nullptr}; }; } // namespace rocshmem diff --git a/projects/rocshmem/src/ipc/context_ipc_device_coll.cpp b/projects/rocshmem/src/ipc/context_ipc_device_coll.cpp new file mode 100644 index 0000000000..f128d76d3f --- /dev/null +++ b/projects/rocshmem/src/ipc/context_ipc_device_coll.cpp @@ -0,0 +1,137 @@ +/****************************************************************************** + * 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 +#if defined(__gfx90a__) + __threadfence_system(); +#endif /* __gfx90a__ */ + 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); +#if defined(__gfx90a__) + __threadfence_system(); +#endif /* __gfx90a__ */ + 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) { +#ifdef USE_COOPERATIVE_GROUPS + cg::grid_group grid = cg::this_grid(); + grid.sync(); + if (0 == grid.thread_rank()) +#else + notifier_->sync(); + if (0 == get_flat_id()) +#endif /* USE_COOPERATIVE_GROUPS */ + { + 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(); +#ifdef USE_COOPERATIVE_GROUPS + grid.sync(); +#else + notifier_->sync(); +#endif /* USE_COOPERATIVE_GROUPS */ +} + +__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/projects/rocshmem/src/ipc/context_ipc_tmpl_device.hpp b/projects/rocshmem/src/ipc/context_ipc_tmpl_device.hpp index 94ef855736..91bdbd45e7 100644 --- a/projects/rocshmem/src/ipc/context_ipc_tmpl_device.hpp +++ b/projects/rocshmem/src/ipc/context_ipc_tmpl_device.hpp @@ -27,6 +27,7 @@ #include "roc_shmem/roc_shmem.hpp" #include "context_ipc_device.hpp" #include "../util.hpp" +#include "ipc_team.hpp" namespace rocshmem { @@ -165,11 +166,46 @@ __device__ void IPCContext::to_all(T *dest, const T *source, int nreduce, 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) { - //broadcast(dst, src, nelems, pe_root_world, pe_start, log_pe_stride, - // pe_size, p_sync); + 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 @@ -177,16 +213,82 @@ __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 @@ -238,58 +340,6 @@ __device__ void IPCContext::get_nbi_wave(T *dest, const T *source, getmem_nbi_wave(dest, source, nelems * sizeof(T), pe); } - -//Wait/test functions -template -__device__ void wait_until(T* ptr, roc_shmem_cmps cmp, T val) { -} - -template -__device__ void wait_until_all(T* ptr, size_t nelems, - const int *status, - roc_shmem_cmps cmp, T val) { -} - -template -__device__ size_t wait_until_any(T* ptr, size_t nelems, - const int *status, - roc_shmem_cmps cmp, T val) { - return 0; -} - -template -__device__ size_t wait_until_some(T* ptr, size_t nelems, - size_t* indices, - const int *status, - roc_shmem_cmps cmp, T val){ - return 0; -} - -template -__device__ void wait_until_all_vector(T* ptr, size_t nelems, - const int *status, - roc_shmem_cmps cmp, T* vals) { -} - -template -__device__ size_t wait_until_any_vector(T* ptr, size_t nelems, - const int *status, - roc_shmem_cmps cmp, T* vals){ - return 0; -} - -template -__device__ size_t wait_until_some_vector(T* ptr, size_t nelems, - size_t* indices, - const int *status, - roc_shmem_cmps cmp, T* vals) { -} - -template -__device__ int test(T* ptr, roc_shmem_cmps cmp, T val) { - return 0; -} - } // namespace rocshmem #endif // LIBRARY_SRC_IPC_CONTEXT_TMPL_DEVICE_HPP_ diff --git a/projects/rocshmem/src/ipc/ipc_context_proxy.hpp b/projects/rocshmem/src/ipc/ipc_context_proxy.hpp index 867f199094..87ca4371e0 100644 --- a/projects/rocshmem/src/ipc/ipc_context_proxy.hpp +++ b/projects/rocshmem/src/ipc/ipc_context_proxy.hpp @@ -41,10 +41,11 @@ class IPCDefaultContextProxy { /* * Placement new the memory which is allocated by proxy_ */ - explicit IPCDefaultContextProxy(IPCBackend* backend) : constructed_{true} { + explicit IPCDefaultContextProxy(IPCBackend* backend, TeamInfo *tinfo) + : constructed_{true} { auto ctx{proxy_.get()}; new (ctx) IPCContext(reinterpret_cast(backend)); - roc_shmem_ctx_t local{ctx, nullptr}; + roc_shmem_ctx_t local{ctx, tinfo}; set_internal_ctx(&local); } diff --git a/projects/rocshmem/src/ipc_policy.hpp b/projects/rocshmem/src/ipc_policy.hpp index c0190198ca..1b84c52613 100644 --- a/projects/rocshmem/src/ipc_policy.hpp +++ b/projects/rocshmem/src/ipc_policy.hpp @@ -64,38 +64,38 @@ class IpcOnImpl { __device__ void ipcCopy_wave(void *dst, void *src, size_t size); - __device__ void ipcFence() { __threadfence(); } + __device__ void ipcFence() { __threadfence_system(); } template __device__ T ipcAMOFetchAdd(T *val, T value) { - return __hip_atomic_fetch_add(val, value, __ATOMIC_RELAXED, - __HIP_MEMORY_SCOPE_AGENT); + return __hip_atomic_fetch_add(val, value, __ATOMIC_SEQ_CST, + __HIP_MEMORY_SCOPE_SYSTEM); } template __device__ T ipcAMOFetchCas(T *val, T cond, T value) { - __hip_atomic_compare_exchange_strong(val, &cond, value, __ATOMIC_RELAXED, - __ATOMIC_RELAXED, - __HIP_MEMORY_SCOPE_AGENT); + __hip_atomic_compare_exchange_strong(val, &cond, value, __ATOMIC_SEQ_CST, + __ATOMIC_SEQ_CST, + __HIP_MEMORY_SCOPE_SYSTEM); return cond; } template __device__ void ipcAMOAdd(T *val, T value) { - __hip_atomic_fetch_add(val, value, __ATOMIC_RELAXED, - __HIP_MEMORY_SCOPE_AGENT); + __hip_atomic_fetch_add(val, value, __ATOMIC_SEQ_CST, + __HIP_MEMORY_SCOPE_SYSTEM); } template __device__ void ipcAMOCas(T *val, T cond, T value) { - __hip_atomic_compare_exchange_strong(val, &cond, value, __ATOMIC_RELAXED, - __ATOMIC_RELAXED, - __HIP_MEMORY_SCOPE_AGENT); + __hip_atomic_compare_exchange_strong(val, &cond, value, __ATOMIC_SEQ_CST, + __ATOMIC_SEQ_CST, + __HIP_MEMORY_SCOPE_SYSTEM); } template __device__ void ipcAMOSet(T *val, T value) { - __hip_atomic_store(val, value, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); + __hip_atomic_store(val, value, __ATOMIC_SEQ_CST, __HIP_MEMORY_SCOPE_SYSTEM); } __device__ void zero_byte_read(int pe) { diff --git a/projects/rocshmem/src/memory/heap_type.hpp b/projects/rocshmem/src/memory/heap_type.hpp index 0008317498..3f48edd11e 100644 --- a/projects/rocshmem/src/memory/heap_type.hpp +++ b/projects/rocshmem/src/memory/heap_type.hpp @@ -40,7 +40,7 @@ namespace rocshmem { #if defined USE_MANAGED_HEAP using HEAP_T = HeapMemory; -#elif defined USE_COHERENT_HEAP || defined USE_CACHED_HEAP +#elif defined USE_COHERENT_HEAP using HEAP_T = HeapMemory; #elif defined USE_HOST_HEAP using HEAP_T = HeapMemory; diff --git a/projects/rocshmem/src/memory/notifier.hpp b/projects/rocshmem/src/memory/notifier.hpp index 0d8cd92f2f..946a9c291e 100644 --- a/projects/rocshmem/src/memory/notifier.hpp +++ b/projects/rocshmem/src/memory/notifier.hpp @@ -20,59 +20,96 @@ * IN THE SOFTWARE. *****************************************************************************/ -/** - * @file notifier.hpp - * - * @brief Contains the notification memory space for threads to communicate - * results with one another. - * - * Assume one thread does work on behalf of other threads (as a leader) and - * that work needs to be communicated to the other threads as a result. - * To expose the result, the threads need to share a memory space where the - * result can be written. The leader thread writes the result out to this - * memory space and all threads synchronize on it. - * - * This class allows the leader thread to notify other threads of the update. - */ - #ifndef LIBRARY_SRC_MEMORY_NOTIFIER_HPP_ #define LIBRARY_SRC_MEMORY_NOTIFIER_HPP_ #include "../device_proxy.hpp" #include "../util.hpp" +#include "../atomic.hpp" namespace rocshmem { +template class Notifier { - public: - __device__ uint64_t read() { return value_; } - __device__ void write(uint64_t val) { - if (is_thread_zero_in_block()) { - value_ = val; - } - publish(); + public: + __device__ uint64_t load() { + return detail::atomic::load(&value_, orders_); } - __device__ void done() { __syncthreads(); } + __device__ void store(uint64_t val) { + detail::atomic::store(&value_, val, orders_); + } - private: - __device__ void publish() { - if (is_thread_zero_in_block()) { - __threadfence(); + __device__ void fence() { + detail::atomic::threadfence(); + } + + __device__ void sync() { + if constexpr (scope == detail::atomic::memory_scope_thread || + scope == detail::atomic::memory_scope_wavefront) { + return; + } + if constexpr (scope == detail::atomic::memory_scope_workgroup) { + __syncthreads(); + return; + } + if constexpr (scope == detail::atomic::memory_scope_system) { + assert(false); + return; + } + + uint32_t done {signal_ + 1}; + __syncthreads(); + + uint32_t retval {0}; + bool executor {!threadIdx.x && !threadIdx.y && !threadIdx.z}; + if (executor) { + retval = detail::atomic::fetch_add(&count_, 1, orders_); + fence(); + } + __syncthreads(); + + if (retval == ((gridDim.x * gridDim.y * gridDim.z) - 1)) { + if (executor) { + detail::atomic::store(&count_, 0, orders_); + fence(); + detail::atomic::fetch_add(&signal_, 1, orders_); + } + } + + if (executor) { + while (detail::atomic::load(&signal_, orders_) != done) { + ; + } } __syncthreads(); } + private: + detail::atomic::rocshmem_memory_orders orders_{}; + uint64_t value_{}; + + uint32_t signal_ {}; + + uint32_t count_ {}; }; -template +template class NotifierProxy { - using ProxyT = DeviceProxy; + using ProxyT = DeviceProxy>; public: - __host__ __device__ Notifier* get() { return proxy_.get(); } + NotifierProxy() { + new (proxy_.get()) Notifier(); + } + + ~NotifierProxy() { + proxy_.get()->~Notifier(); + } + + __host__ __device__ Notifier* get() { return proxy_.get(); } private: ProxyT proxy_{}; diff --git a/projects/rocshmem/src/memory/slab_heap.cpp b/projects/rocshmem/src/memory/slab_heap.cpp index 89067b78da..faf3a84b1b 100644 --- a/projects/rocshmem/src/memory/slab_heap.cpp +++ b/projects/rocshmem/src/memory/slab_heap.cpp @@ -75,9 +75,12 @@ __device__ void SlabHeap::malloc(void** ptr, size_t size) { * Notify other threads in block about the allocation result. */ auto notifier{notifier_.get()}; - notifier->write(ptr_deref_u64); - uint64_t notification_u64{notifier->read()}; - notifier->done(); + if (!threadIdx.x) { + notifier->store(ptr_deref_u64); + notifier->fence(); + } + __syncthreads(); + uint64_t notification_u64{notifier->load()}; /* * Write to the ptr parameter (to return it back up the call stack). diff --git a/projects/rocshmem/src/memory/slab_heap.hpp b/projects/rocshmem/src/memory/slab_heap.hpp index a3655b27c1..171332bee3 100644 --- a/projects/rocshmem/src/memory/slab_heap.hpp +++ b/projects/rocshmem/src/memory/slab_heap.hpp @@ -48,7 +48,7 @@ class SlabHeap { /** * @brief Helper type for notifier */ - using NOTIFIER_PROXY_T = NotifierProxy; + using NOTIFIER_PROXY_T = NotifierProxy; /** * @brief Helper type for notifier diff --git a/projects/rocshmem/src/reverse_offload/backend_ro.cpp b/projects/rocshmem/src/reverse_offload/backend_ro.cpp index 1c1bf4645c..96471181e3 100644 --- a/projects/rocshmem/src/reverse_offload/backend_ro.cpp +++ b/projects/rocshmem/src/reverse_offload/backend_ro.cpp @@ -94,7 +94,9 @@ ROBackend::ROBackend(MPI_Comm comm) default_block_handle_proxy_ = DefaultBlockHandleProxyT( bp->g_ret, bp->atomic_ret, &queue_, &ipcImpl, hdp_proxy_.get()); - default_context_proxy_ = DefaultContextProxyT(this); + + TeamInfo *tinfo = team_tracker.get_team_world()->tinfo_wrt_world; + default_context_proxy_ = DefaultContextProxyT(this, tinfo); block_handle_proxy_ = BlockHandleProxyT(bp->g_ret, bp->atomic_ret, &queue_, &ipcImpl, hdp_proxy_.get()); diff --git a/projects/rocshmem/src/reverse_offload/context_proxy.hpp b/projects/rocshmem/src/reverse_offload/context_proxy.hpp index d89281d921..4ae94d3851 100644 --- a/projects/rocshmem/src/reverse_offload/context_proxy.hpp +++ b/projects/rocshmem/src/reverse_offload/context_proxy.hpp @@ -42,10 +42,11 @@ class DefaultContextProxy { /* * Placement new the memory which is allocated by proxy_ */ - explicit DefaultContextProxy(ROBackend* backend) : constructed_{true} { + explicit DefaultContextProxy(ROBackend* backend, TeamInfo *tinfo) + : constructed_{true} { auto ctx{proxy_.get()}; new (ctx) ROContext(reinterpret_cast(backend), -1); - roc_shmem_ctx_t local{ctx, nullptr}; + roc_shmem_ctx_t local{ctx, tinfo}; set_internal_ctx(&local); } diff --git a/projects/rocshmem/src/reverse_offload/context_ro_device.cpp b/projects/rocshmem/src/reverse_offload/context_ro_device.cpp index 6182ce0917..86e12c7c38 100644 --- a/projects/rocshmem/src/reverse_offload/context_ro_device.cpp +++ b/projects/rocshmem/src/reverse_offload/context_ro_device.cpp @@ -50,6 +50,9 @@ __host__ ROContext::ROContext(Backend *b, size_t block_id) block_handle = &block_base[block_id]; } ro_net_win_id = block_id % backend->ro_window_proxy_->MAX_NUM_WINDOWS; + + ipcImpl_.ipc_bases = b->ipcImpl.ipc_bases; + ipcImpl_.shm_size = b->ipcImpl.shm_size; } __device__ void ROContext::putmem(void *dest, const void *source, size_t nelems, @@ -389,24 +392,6 @@ __device__ bool enough_space(BlockHandle *h, uint64_t required) { return (h->queue_size - (h->write_index - h->read_index)) >= required; } -__device__ void refresh_volatile_dwordx2(volatile uint64_t *assigned_value, - volatile uint64_t *read_value) { - __asm__ volatile( - "global_load_dwordx2 %0 %1 off glc slc\n " - "s_waitcnt vmcnt(0)" - : "=v"(*assigned_value) - : "v"(read_value)); -} - -__device__ void refresh_volatile_sbyte(volatile int *assigned_value, - volatile char *read_value) { - __asm__ volatile( - "global_load_sbyte %0 %1 off glc slc\n " - "s_waitcnt vmcnt(0)" - : "=v"(*assigned_value) - : "v"(read_value)); -} - __device__ void acquire_lock(BlockHandle *handle) { while(atomicCAS((uint64_t *)&handle->lock, 0, 1) == 1) ; } diff --git a/projects/rocshmem/src/roc_shmem_gpu.cpp b/projects/rocshmem/src/roc_shmem_gpu.cpp index 086a6deed1..1afd5b10ff 100644 --- a/projects/rocshmem/src/roc_shmem_gpu.cpp +++ b/projects/rocshmem/src/roc_shmem_gpu.cpp @@ -262,9 +262,9 @@ __device__ int roc_shmem_wg_ctx_create(long option, roc_shmem_ctx_t *ctx) { GPU_DPRINTF("Function: roc_shmem_ctx_create\n"); bool result{true}; if (get_flat_block_id() == 0) { + ctx->team_opaque = reinterpret_cast(ROC_SHMEM_CTX_DEFAULT.team_opaque); device_backend_proxy->create_ctx(option, ctx); reinterpret_cast(ctx->ctx_opaque)->setFence(option); - ctx->team_opaque = nullptr; } __syncthreads(); return result == true ? 0 : -1; @@ -279,11 +279,11 @@ __device__ int roc_shmem_wg_team_create_ctx(roc_shmem_team_t team, long options, bool result{true}; if (get_flat_block_id() == 0) { - result = device_backend_proxy->create_ctx(options, ctx); - reinterpret_cast(ctx->ctx_opaque)->setFence(options); Team *team_obj{get_internal_team(team)}; TeamInfo *info_wrt_world = team_obj->tinfo_wrt_world; ctx->team_opaque = info_wrt_world; + result = device_backend_proxy->create_ctx(options, ctx); + reinterpret_cast(ctx->ctx_opaque)->setFence(options); } __syncthreads(); @@ -412,7 +412,9 @@ __device__ void roc_shmem_ctx_fence(roc_shmem_ctx_t ctx) { __device__ void roc_shmem_ctx_fence(roc_shmem_ctx_t ctx, int pe) { GPU_DPRINTF("Function: roc_shmem_ctx_fence\n"); - get_internal_ctx(ctx)->fence(pe); + int pe_in_world = translate_pe(ctx, pe); + + get_internal_ctx(ctx)->fence(pe_in_world); } __device__ void roc_shmem_ctx_quiet(roc_shmem_ctx_t ctx) { diff --git a/projects/rocshmem/src/util.cpp b/projects/rocshmem/src/util.cpp index 11122d7b16..7318b14208 100644 --- a/projects/rocshmem/src/util.cpp +++ b/projects/rocshmem/src/util.cpp @@ -40,15 +40,6 @@ typedef struct device_agent { std::vector gpu_agents; std::vector cpu_agents; -__device__ uint64_t __read_clock() { - uint64_t clock{}; - asm volatile( - "s_memrealtime %0\n" - "s_waitcnt lgkmcnt(0)\n" - : "=s"(clock)); - return clock; -} - hsa_status_t rocm_hsa_amd_memory_pool_callback( hsa_amd_memory_pool_t memory_pool, void* data) { hsa_amd_memory_pool_global_flag_t pool_flag{}; diff --git a/projects/rocshmem/src/util.hpp b/projects/rocshmem/src/util.hpp index 67fdac6786..c02f891dc4 100644 --- a/projects/rocshmem/src/util.hpp +++ b/projects/rocshmem/src/util.hpp @@ -29,60 +29,12 @@ #include +#include "assembly.hpp" #include "config.h" // NOLINT(build/include_subdir) #include "constants.hpp" namespace rocshmem { -#define DO_PRAGMA(x) _Pragma(#x) -#define NOWARN(warnoption, ...) \ - DO_PRAGMA(GCC diagnostic push) \ - DO_PRAGMA(GCC diagnostic ignored #warnoption) \ - __VA_ARGS__ \ - DO_PRAGMA(GCC diagnostic pop) - -__device__ __forceinline__ int uncached_load_ubyte(uint8_t* src) { - int ret; - __asm__ volatile( - "global_load_ubyte %0 %1 off glc slc \n" - "s_waitcnt vmcnt(0)" - : "=v"(ret) - : "v"(src)); - return ret; -} - -/* Ignore the warning about deprecated volatile. - * The only usage of volatile is to force the compiler to generate - * the assembly instruction. If volatile is omitted, the compiler - * will NOT generate the non-temporal load or the waitcnt. - */ -// clang-format off -NOWARN(-Wdeprecated-volatile, - template __device__ __forceinline__ T uncached_load(T* src) { - T ret; - switch (sizeof(T)) { - case 4: - __asm__ volatile( - "global_load_dword %0 %1 off glc slc \n" - "s_waitcnt vmcnt(0)" - : "=v"(ret) - : "v"(src)); - break; - case 8: - __asm__ volatile( - "global_load_dwordx2 %0 %1 off glc slc \n" - "s_waitcnt vmcnt(0)" - : "=v"(ret) - : "v"(src)); - break; - default: - break; - } - return ret; - } -) -// clang-format on - #define LOAD(VAR) __atomic_load_n((VAR), __ATOMIC_SEQ_CST) #define STORE(DST, SRC) __atomic_store_n((DST), (SRC), __ATOMIC_SEQ_CST) @@ -96,8 +48,6 @@ NOWARN(-Wdeprecated-volatile, } \ } -#define SFENCE() asm volatile("sfence" ::: "memory") - #ifdef DEBUG #define DPRINTF(...) \ do { \ @@ -123,19 +73,6 @@ NOWARN(-Wdeprecated-volatile, extern const int gpu_clock_freq_mhz; /* Device-side internal functions */ -__device__ __forceinline__ void __roc_inv() { asm volatile("buffer_wbinvl1;"); } - -__device__ __forceinline__ void __roc_flush() { -#if defined USE_COHERENT_HEAP || defined USE_CACHED_HEAP -#if __gfx90a__ - asm volatile("s_dcache_wb;"); - asm volatile("buffer_wbl2;"); -#endif -#endif -} - -__device__ uint64_t __read_clock(); - __device__ __forceinline__ uint32_t lowerID() { return __ffsll(__ballot(1)) - 1; } @@ -156,6 +93,7 @@ __device__ __forceinline__ bool is_thread_zero_in_block() { __device__ __forceinline__ bool is_block_zero_in_grid() { return hipBlockIdx_x == 0 && hipBlockIdx_y == 0 && hipBlockIdx_z == 0; } + /* * Returns the number of threads in the caller's flattened thread block. */ @@ -163,6 +101,13 @@ __device__ __forceinline__ int get_flat_block_size() { return hipBlockDim_x * hipBlockDim_y * hipBlockDim_z; } +/* + * Returns the number of threads in the caller's flattened grid. + */ +__device__ __forceinline__ int get_flat_grid_size() { + return get_flat_block_size() * hipGridDim_x * hipGridDim_y * hipGridDim_z; +} + /* * Returns the flattened thread index of the calling thread within its * thread block. @@ -212,7 +157,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...); @@ -222,29 +167,6 @@ __device__ void gpu_dprintf(const char* fmt, const Args&... args) { } } -__device__ __forceinline__ void store_asm(uint8_t* val, uint8_t* dst, - int size) { - switch (size) { - case 2: { - int16_t val16{*(reinterpret_cast(val))}; - asm volatile("flat_store_short %0 %1 glc slc" : : "v"(dst), "v"(val16)); - break; - } - case 4: { - int32_t val32{*(reinterpret_cast(val))}; - asm volatile("flat_store_dword %0 %1 glc slc" : : "v"(dst), "v"(val32)); - break; - } - case 8: { - int64_t val64{*(reinterpret_cast(val))}; - asm volatile("flat_store_dwordx2 %0 %1 glc slc" : : "v"(dst), "v"(val64)); - break; - } - default: - break; - } -} - __device__ __forceinline__ void memcpy(void* dst, void* src, size_t size) { uint8_t* dst_bytes{static_cast(dst)}; uint8_t* src_bytes{static_cast(src)}; diff --git a/projects/rocshmem/tests/functional_tests/sync_tester.cpp b/projects/rocshmem/tests/functional_tests/sync_tester.cpp index 54d4a99f99..1afecaa45c 100644 --- a/projects/rocshmem/tests/functional_tests/sync_tester.cpp +++ b/projects/rocshmem/tests/functional_tests/sync_tester.cpp @@ -83,9 +83,21 @@ void SyncTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop, roc_shmem_team_split_strided(ROC_SHMEM_TEAM_WORLD, 0, 1, n_pes, nullptr, 0, &team_sync_world_dup); +#ifdef USE_COOPERATIVE_GROUPS + void* kernelParams[] = {(void*)&loop, + (void*)&args.skip, + (void*)&timer, + (void*)&_type, + (void*)&_shmem_context, + (void*)&team_sync_world_dup}; + + CHECK_HIP(hipLaunchCooperativeKernel(SyncTest, gridSize, blockSize, + kernelParams, shared_bytes, stream)); +#else hipLaunchKernelGGL(SyncTest, gridSize, blockSize, shared_bytes, stream, loop, args.skip, timer, _type, _shmem_context, team_sync_world_dup); +#endif /* USE_COOPERATIVE_GROUPS */ num_msgs = (loop + args.skip) * gridSize.x; num_timed_msgs = loop; diff --git a/projects/rocshmem/tests/unit_tests/CMakeLists.txt b/projects/rocshmem/tests/unit_tests/CMakeLists.txt index 349bb2c2bf..b47da1c154 100644 --- a/projects/rocshmem/tests/unit_tests/CMakeLists.txt +++ b/projects/rocshmem/tests/unit_tests/CMakeLists.txt @@ -90,6 +90,7 @@ target_sources( free_list_gtest.cpp #context_ipc_gtest.cpp ipc_impl_simple_coarse_gtest.cpp + ipc_impl_simple_fine_gtest.cpp ) ############################################################################### diff --git a/projects/rocshmem/tests/unit_tests/ipc_impl_simple_coarse_gtest.cpp b/projects/rocshmem/tests/unit_tests/ipc_impl_simple_coarse_gtest.cpp index dbf72923b5..37833f9cfe 100644 --- a/projects/rocshmem/tests/unit_tests/ipc_impl_simple_coarse_gtest.cpp +++ b/projects/rocshmem/tests/unit_tests/ipc_impl_simple_coarse_gtest.cpp @@ -33,6 +33,7 @@ TEST_F(IPCImplSimpleCoarseTestFixture, MPI_num_pes) { } TEST_F(IPCImplSimpleCoarseTestFixture, IPC_bases) { + ASSERT_NE(ipc_impl_.ipc_bases, nullptr); for(int i{0}; i < mpi_.num_pes(); i++) { ASSERT_NE(ipc_impl_.ipc_bases[i], nullptr); } diff --git a/projects/rocshmem/tests/unit_tests/ipc_impl_simple_coarse_gtest.hpp b/projects/rocshmem/tests/unit_tests/ipc_impl_simple_coarse_gtest.hpp index 02dfd8c55a..083d73f418 100644 --- a/projects/rocshmem/tests/unit_tests/ipc_impl_simple_coarse_gtest.hpp +++ b/projects/rocshmem/tests/unit_tests/ipc_impl_simple_coarse_gtest.hpp @@ -218,8 +218,6 @@ class IPCImplSimpleCoarseTestFixture : public ::testing::Test { protected: std::vector golden_; - std::vector output_; - HEAP_T heap_mem_ {}; MPI_T mpi_ {heap_mem_.get_ptr(), heap_mem_.get_size()}; diff --git a/projects/rocshmem/tests/unit_tests/ipc_impl_simple_fine_gtest.cpp b/projects/rocshmem/tests/unit_tests/ipc_impl_simple_fine_gtest.cpp new file mode 100644 index 0000000000..71cb85dc9e --- /dev/null +++ b/projects/rocshmem/tests/unit_tests/ipc_impl_simple_fine_gtest.cpp @@ -0,0 +1,1037 @@ +/****************************************************************************** + * 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_fine_gtest.hpp" + +using namespace rocshmem; + +TEST_F(IPCImplSimpleFineTestFixture, ptr_check) { + ASSERT_NE(heap_mem_.get_ptr(), nullptr); +} + +TEST_F(IPCImplSimpleFineTestFixture, MPI_num_pes) { + ASSERT_EQ(mpi_.num_pes(), 2); +} + +TEST_F(IPCImplSimpleFineTestFixture, IPC_bases) { + ASSERT_NE(ipc_impl_.ipc_bases, nullptr); + for(int i{0}; i < mpi_.num_pes(); i++) { + ASSERT_NE(ipc_impl_.ipc_bases[i], nullptr); + } +} + +TEST_F(IPCImplSimpleFineTestFixture, golden_1048576_int) { + iota_golden(1048576); + validate_golden(1048576); +} + +//============================================================================= + +TEST_F(IPCImplSimpleFineTestFixture, write_wg_1x1x1_1x1x1_32_int) { + dim3 grid {1,1,1}; + dim3 block {1,1,1}; + write_wg(grid, block, 32); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wg_1x1x1_1024x1x1_32_int) { + dim3 grid {1,1,1}; + dim3 block {1024,1,1}; + write_wg(grid, block, 32); +} + +//============================================================================= + +TEST_F(IPCImplSimpleFineTestFixture, read_wg_1x1x1_1024x1x1_32_int) { + dim3 grid {1,1,1}; + dim3 block {1024,1,1}; + read_wg(grid, block, 32); +} + +//============================================================================= + +TEST_F(IPCImplSimpleFineTestFixture, write_wg_1x1x1_1x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {1,1,1}; + write_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wg_1x1x1_2x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {2,1,1}; + write_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wg_1x1x1_4x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {4,1,1}; + write_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wg_1x1x1_8x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {8,1,1}; + write_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wg_1x1x1_16x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {16,1,1}; + write_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wg_1x1x1_32x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {32,1,1}; + write_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wg_1x1x1_64x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {64,1,1}; + write_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wg_1x1x1_128x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {128,1,1}; + write_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wg_1x1x1_256x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {256,1,1}; + write_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wg_1x1x1_512x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {512,1,1}; + write_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wg_1x1x1_768x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {768,1,1}; + write_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wg_1x1x1_1024x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {1024,1,1}; + write_wg(grid, block, 1048576); +} + +//============================================================================= + +TEST_F(IPCImplSimpleFineTestFixture, read_wg_1x1x1_1x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {1,1,1}; + read_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wg_1x1x1_2x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {2,1,1}; + read_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wg_1x1x1_4x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {4,1,1}; + read_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wg_1x1x1_8x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {8,1,1}; + read_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wg_1x1x1_16x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {16,1,1}; + read_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wg_1x1x1_32x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {32,1,1}; + read_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wg_1x1x1_64x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {64,1,1}; + read_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wg_1x1x1_128x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {128,1,1}; + read_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wg_1x1x1_256x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {256,1,1}; + read_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wg_1x1x1_512x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {512,1,1}; + read_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wg_1x1x1_768x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {768,1,1}; + read_wg(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wg_1x1x1_1024x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {1024,1,1}; + read_wg(grid, block, 1048576); +} + +//============================================================================= + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_64x1x1_1_int) { + dim3 grid {1,1,1}; + dim3 block {64,1,1}; + write_wave(grid, block, 1); +} + +//============================================================================= + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_64x1x1_1_int) { + dim3 grid {1,1,1}; + dim3 block {64,1,1}; + read_wave(grid, block, 1); +} + +//============================================================================= + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_64x1x1_32_int) { + dim3 grid {1,1,1}; + dim3 block {64,1,1}; + write_wave(grid, block, 32); +} + +//============================================================================= + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_64x1x1_32_int) { + dim3 grid {1,1,1}; + dim3 block {64,1,1}; + read_wave(grid, block, 32); +} + +//============================================================================= + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_1x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {1,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_2x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {2,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_3x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {3,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_4x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {4,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_5x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {5,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_6x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {6,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_7x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {7,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_8x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {8,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_9x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {9,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_10x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {10,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_11x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {11,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_12x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {12,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_13x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {13,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_14x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {14,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_15x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {15,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_16x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {16,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_17x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {17,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_18x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {18,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_19x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {19,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_20x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {20,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_21x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {21,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_22x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {22,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_23x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {23,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_24x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {24,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_25x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {25,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_26x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {26,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_27x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {27,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_28x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {28,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_29x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {29,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_30x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {30,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_31x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {31,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_32x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {32,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_33x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {33,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_34x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {34,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_35x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {35,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_36x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {36,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_37x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {37,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_38x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {38,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_39x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {39,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_40x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {40,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_41x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {41,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_42x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {42,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_43x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {43,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_44x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {44,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_45x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {45,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_46x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {46,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_47x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {47,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_48x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {48,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_49x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {49,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_50x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {50,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_51x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {51,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_52x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {52,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_53x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {53,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_54x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {54,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_55x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {55,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_56x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {56,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_57x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {57,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_58x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {58,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_59x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {59,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_60x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {60,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_61x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {61,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_62x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {62,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_63x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {63,1,1}; + write_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, write_wave_1x1x1_64x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {64,1,1}; + write_wave(grid, block, 1048576); +} + +//============================================================================= + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_1x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {1,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_2x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {2,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_3x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {3,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_4x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {4,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_5x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {5,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_6x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {6,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_7x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {7,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_8x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {8,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_9x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {9,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_10x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {10,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_11x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {11,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_12x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {12,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_13x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {13,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_14x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {14,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_15x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {15,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_16x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {16,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_17x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {17,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_18x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {18,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_19x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {19,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_20x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {20,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_21x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {21,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_22x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {22,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_23x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {23,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_24x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {24,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_25x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {25,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_26x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {26,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_27x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {27,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_28x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {28,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_29x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {29,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_30x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {30,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_31x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {31,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_32x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {32,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_33x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {33,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_34x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {34,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_35x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {35,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_36x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {36,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_37x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {37,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_38x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {38,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_39x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {39,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_40x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {40,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_41x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {41,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_42x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {42,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_43x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {43,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_44x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {44,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_45x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {45,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_46x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {46,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_47x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {47,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_48x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {48,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_49x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {49,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_50x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {50,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_51x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {51,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_52x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {52,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_53x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {53,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_54x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {54,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_55x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {55,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_56x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {56,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_57x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {57,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_58x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {58,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_59x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {59,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_60x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {60,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_61x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {61,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_62x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {62,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_63x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {63,1,1}; + read_wave(grid, block, 1048576); +} + +TEST_F(IPCImplSimpleFineTestFixture, read_wave_1x1x1_64x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {64,1,1}; + read_wave(grid, block, 1048576); +} + +//============================================================================= + +TEST_F(IPCImplSimpleFineTestFixture, write_1x1x1_1x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {1,1,1}; + write(grid, block, 1048576); +} + +//============================================================================= + +TEST_F(IPCImplSimpleFineTestFixture, read_1x1x1_1x1x1_1048576_int) { + dim3 grid {1,1,1}; + dim3 block {1,1,1}; + read(grid, block, 1048576); +} + + diff --git a/projects/rocshmem/tests/unit_tests/ipc_impl_simple_fine_gtest.hpp b/projects/rocshmem/tests/unit_tests/ipc_impl_simple_fine_gtest.hpp new file mode 100644 index 0000000000..ba4c20e288 --- /dev/null +++ b/projects/rocshmem/tests/unit_tests/ipc_impl_simple_fine_gtest.hpp @@ -0,0 +1,344 @@ +/****************************************************************************** + * 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_FINE_GTEST_HPP +#define ROCSHMEM_IPC_IMPL_SIMPLE_FINE_GTEST_HPP + +#include "gtest/gtest.h" + +#include +#include + +#include "../src/atomic.hpp" +#include "../src/ipc_policy.hpp" +#include "../src/memory/notifier.hpp" +#include "../src/memory/symmetric_heap.hpp" +#include "../src/util.hpp" + +namespace rocshmem { + +enum TestType { + READ = 0, + WRITE = 1 +}; + +const uint32_t SIGNAL_OFFSET {67108864}; + +__device__ +void +validator(bool *error, int *golden, int *dest, size_t bytes) { + size_t elements {bytes / sizeof(int)}; + for (int i {get_flat_id()}; i < elements; i += get_flat_grid_size()) { + if (golden[i] != dest[i]) { + printf("golden[%d] %d != dest[%d] %d\n", i, golden[i], i, dest[i]); + *error = true; + } + } +} + +template +__global__ +void +kernel_put_with_signal_validator(bool *error, int *golden, int *dest, size_t bytes, NotifierT *notifier) { + detail::atomic::rocshmem_memory_orders orders{}; + if (!get_flat_id()) { + while (detail::atomic::load(dest + SIGNAL_OFFSET, orders) == 0) { + ; + } + } + notifier->sync(); + validator(error, golden, dest, bytes); +} + +template +__global__ +void +kernel_simple_fine_copy(IpcImpl *ipc_impl, bool *error, int *golden, int *src, int *dest, size_t bytes, TestType test, NotifierT *notifier) { + if (!get_flat_id()) { + ipc_impl->ipcCopy(dest, src, bytes); + ipc_impl->ipcFence(); + if (test == WRITE) { + ipc_impl->ipcAMOFetchAdd(dest + SIGNAL_OFFSET, 1); + } + } + if (test == READ) { + notifier->sync(); + validator(error, golden, dest, bytes); + } +} + +template +__global__ +void +kernel_simple_fine_copy_wg(IpcImpl *ipc_impl, bool *error, int *golden, int *src, int *dest, size_t bytes, TestType test, NotifierT *notifier) { + if (!blockIdx.x) { + ipc_impl->ipcCopy_wg(dest, src, bytes); + ipc_impl->ipcFence(); + if (test == WRITE) { + if (!threadIdx.x) { + ipc_impl->ipcAMOFetchAdd(dest + SIGNAL_OFFSET, 1); + } + } + } + if (test == READ) { + notifier->sync(); + validator(error, golden, dest, bytes); + } +} + +template +__global__ +void +kernel_simple_fine_copy_wave(IpcImpl *ipc_impl, bool *error, int *golden, int *src, int *dest, size_t bytes, TestType test, NotifierT *notifier) { + if (!blockIdx.x && threadIdx.x < 64) { + ipc_impl->ipcCopy_wave(dest, src, bytes); + ipc_impl->ipcFence(); + if (test == WRITE) { + if (!threadIdx.x) { + ipc_impl->ipcAMOFetchAdd(dest + SIGNAL_OFFSET, 1); + } + } + } + __syncthreads(); + if (test == READ) { + notifier->sync(); + validator(error, golden, dest, bytes); + } +} + +class IPCImplSimpleFineTestFixture : public ::testing::Test { + using HEAP_T = HeapMemory; + using MPI_T = RemoteHeapInfo; + using NotifierT = Notifier; + using NotifierProxyT = NotifierProxy; + using FN_T1 = void (*)(IpcImpl*, bool*, int*, int*, int*, size_t, TestType, NotifierT*); + using FN_T2 = void (*)(bool*, int*, int*, size_t, NotifierT*); + + public: + IPCImplSimpleFineTestFixture() { + 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)); + + assert(error_dptr_ == nullptr); + hip_allocator_.allocate((void**)&error_dptr_, sizeof(bool)); + *error_dptr_ = false; + } + + ~IPCImplSimpleFineTestFixture() { + if (ipc_impl_dptr_) { + hip_allocator_.deallocate(ipc_impl_dptr_); + } + if (error_dptr_) { + hip_allocator_.deallocate(error_dptr_); + } + if (golden_dptr_) { + hip_allocator_.deallocate(golden_dptr_); + } + + ipc_impl_.ipcHostStop(); + } + + void launch(FN_T1 f, const dim3 grid, const dim3 block, int* src, int* dest, size_t bytes, TestType test) { + f<<>>(ipc_impl_dptr_, error_dptr_, golden_dptr_, src, dest, bytes, test, notifier_.get()); + CHECK_HIP(hipStreamSynchronize(nullptr)); + } + + void launch(FN_T2 f, const dim3 grid, const dim3 block, int* dest, size_t bytes) { + f<<>>(error_dptr_, golden_dptr_, dest, bytes, notifier_.get()); + CHECK_HIP(hipStreamSynchronize(nullptr)); + } + + void write(const dim3 grid, const dim3 block, size_t elems) { + iota_golden(elems); + initialize_signal(WRITE); + initialize_src_buffer(WRITE); + copy(WRITE, grid, block); + check_device_validation_errors(WRITE); + } + + void write_wg(const dim3 grid, const dim3 block, size_t elems) { + iota_golden(elems); + initialize_signal(WRITE); + initialize_src_buffer(WRITE); + copy_wg(WRITE, grid, block); + check_device_validation_errors(WRITE); + } + + void write_wave(const dim3 grid, const dim3 block, size_t elems) { + iota_golden(elems); + initialize_signal(WRITE); + initialize_src_buffer(WRITE); + copy_wave(WRITE, grid, block); + check_device_validation_errors(WRITE); + } + + void read(const dim3 grid, const dim3 block, size_t elems) { + iota_golden(elems); + initialize_signal(READ); + initialize_src_buffer(READ); + copy(READ, grid, block); + check_device_validation_errors(READ); + } + + void read_wg(const dim3 grid, const dim3 block, size_t elems) { + iota_golden(elems); + initialize_signal(READ); + initialize_src_buffer(READ); + copy_wg(READ, grid, block); + check_device_validation_errors(READ); + } + + void read_wave(const dim3 grid, const dim3 block, size_t elems) { + iota_golden(elems); + initialize_signal(READ); + initialize_src_buffer(READ); + copy_wave(READ, grid, block); + check_device_validation_errors(READ); + } + + void iota_golden(size_t elems) { + golden_.resize(elems); + std::iota(golden_.begin(), golden_.end(), 0); + + assert(golden_dptr_ == nullptr); + size_t golden_dptr_bytes {golden_.size() * sizeof(int)}; + hip_allocator_.allocate((void**)&golden_dptr_, golden_dptr_bytes); + CHECK_HIP(hipMemcpy(golden_dptr_, golden_.data(), golden_dptr_bytes, hipMemcpyHostToDevice)); + } + + 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_signal(TestType test) { + bool is_write_test = test; + if (is_write_test && mpi_.my_pe() == 0) { + int *dest = reinterpret_cast(ipc_impl_.ipc_bases[1]); + *(dest + SIGNAL_OFFSET) = 0; + } + } + + 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)); + } + + 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_T1 fn, const dim3 grid, const dim3 block) { + size_t bytes = golden_.size() * sizeof(int); + if (mpi_.my_pe()) { + mpi_.barrier(); + if (test == WRITE) { + int *dest = reinterpret_cast(ipc_impl_.ipc_bases[1]); + FN_T2 val_fn = kernel_put_with_signal_validator; + launch(val_fn, grid, block, dest, bytes); + } + 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]); + } + mpi_.barrier(); + launch(fn, grid, block, src, dest, bytes, test); + mpi_.barrier(); + } + + void copy(TestType test, dim3 grid, dim3 block) { + execute(test, kernel_simple_fine_copy, grid, block); + } + + void copy_wg(TestType test, dim3 grid, dim3 block) { + execute(test, kernel_simple_fine_copy_wg, grid, block); + } + + void copy_wave(TestType test, dim3 grid, dim3 block) { + execute(test, kernel_simple_fine_copy_wave, grid, block); + } + + void check_device_validation_errors(TestType test) { + if (!pe_validates_dest_buffer(test)) { + return; + } + ASSERT_EQ(*error_dptr_, false); + } + + 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: + HIPDefaultFinegrainedAllocator hip_allocator_ {}; + + NotifierProxyT notifier_ {}; + + HEAP_T heap_mem_ {}; + + MPI_T mpi_ {heap_mem_.get_ptr(), heap_mem_.get_size()}; + + std::vector golden_; + + int *golden_dptr_ {nullptr}; + + IpcImpl ipc_impl_ {}; + + IpcImpl *ipc_impl_dptr_ {nullptr}; + + bool *error_dptr_ {nullptr}; +}; + +} // namespace rocshmem + +#endif // ROCSHMEM_IPC_IMPL_SIMPLE_FINE_GTEST_HPP diff --git a/projects/rocshmem/tests/unit_tests/notifier_gtest.cpp b/projects/rocshmem/tests/unit_tests/notifier_gtest.cpp index 9f79c62397..d3b699e3d9 100644 --- a/projects/rocshmem/tests/unit_tests/notifier_gtest.cpp +++ b/projects/rocshmem/tests/unit_tests/notifier_gtest.cpp @@ -28,30 +28,98 @@ using namespace rocshmem; ******************************* Fixture Tests ******************************* *****************************************************************************/ -TEST_F(NotifierTestFixture, run_all_threads_once_1_1) { +TEST_F(NotifierBlockTestFixture, run_all_threads_once_1_1) { run_all_threads_once(1, 1); } -TEST_F(NotifierTestFixture, run_all_threads_once_2_1) { +TEST_F(NotifierBlockTestFixture, run_all_threads_once_2_1) { run_all_threads_once(2, 1); } -TEST_F(NotifierTestFixture, run_all_threads_once_64_1) { +TEST_F(NotifierBlockTestFixture, run_all_threads_once_64_1) { run_all_threads_once(64, 1); } -TEST_F(NotifierTestFixture, run_all_threads_once_128_1) { +TEST_F(NotifierBlockTestFixture, run_all_threads_once_128_1) { run_all_threads_once(128, 1); } -TEST_F(NotifierTestFixture, run_all_threads_once_256_1) { +TEST_F(NotifierBlockTestFixture, run_all_threads_once_256_1) { run_all_threads_once(256, 1); } -TEST_F(NotifierTestFixture, run_all_threads_once_512_1) { +TEST_F(NotifierBlockTestFixture, run_all_threads_once_512_1) { run_all_threads_once(512, 1); } -TEST_F(NotifierTestFixture, run_all_threads_once_1024_1) { +TEST_F(NotifierBlockTestFixture, run_all_threads_once_1024_1) { run_all_threads_once(1024, 1); } + +TEST_F(NotifierAgentTestFixture, run_all_threads_once_1_1) { + run_all_threads_once(1, 1); +} + +TEST_F(NotifierAgentTestFixture, run_all_threads_once_2_1) { + run_all_threads_once(2, 1); +} + +TEST_F(NotifierAgentTestFixture, run_all_threads_once_64_1) { + run_all_threads_once(64, 1); +} + +TEST_F(NotifierAgentTestFixture, run_all_threads_once_128_1) { + run_all_threads_once(128, 1); +} + +TEST_F(NotifierAgentTestFixture, run_all_threads_once_256_1) { + run_all_threads_once(256, 1); +} + +TEST_F(NotifierAgentTestFixture, run_all_threads_once_512_1) { + run_all_threads_once(512, 1); +} + +TEST_F(NotifierAgentTestFixture, run_all_threads_once_1024_1) { + run_all_threads_once(1024, 1); +} + +TEST_F(NotifierAgentTestFixture, run_all_threads_once_1_2) { + run_all_threads_once(1, 2); +} + +TEST_F(NotifierAgentTestFixture, run_all_threads_once_1024_2) { + run_all_threads_once(1024, 2); +} + +TEST_F(NotifierAgentTestFixture, run_all_threads_once_1_4) { + run_all_threads_once(1, 4); +} + +TEST_F(NotifierAgentTestFixture, run_all_threads_once_1024_4) { + run_all_threads_once(1024, 4); +} + +TEST_F(NotifierAgentTestFixture, run_all_threads_once_1_8) { + run_all_threads_once(1, 8); +} + +TEST_F(NotifierAgentTestFixture, run_all_threads_once_1024_8) { + run_all_threads_once(1024, 8); +} + +TEST_F(NotifierAgentTestFixture, run_all_threads_once_1_32) { + run_all_threads_once(1, 32); +} + +TEST_F(NotifierAgentTestFixture, run_all_threads_once_1024_32) { + run_all_threads_once(1024, 32); +} + +TEST_F(NotifierAgentTestFixture, run_all_threads_once_1_38) { + run_all_threads_once(1, 38); // MI300 CPX +} + +TEST_F(NotifierAgentTestFixture, run_all_threads_once_1024_38) { + run_all_threads_once(1024, 38); // MI300 CPX +} diff --git a/projects/rocshmem/tests/unit_tests/notifier_gtest.hpp b/projects/rocshmem/tests/unit_tests/notifier_gtest.hpp index ecb2e7a619..e130159b2e 100644 --- a/projects/rocshmem/tests/unit_tests/notifier_gtest.hpp +++ b/projects/rocshmem/tests/unit_tests/notifier_gtest.hpp @@ -44,66 +44,45 @@ static const uint64_t NOTIFIER_OFFSET {0x100B00}; inline __device__ void write_to_memory(uint8_t* raw_memory) { - auto thread_idx {get_flat_block_id()}; + auto thread_idx {get_flat_id()}; raw_memory[thread_idx] = THREAD_VALUE; __threadfence(); } +template __global__ void all_threads_once(uint8_t* raw_memory, - Notifier* notifier) { - notifier->write(NOTIFIER_OFFSET); - uint64_t offset_u64 {notifier->read()}; - notifier->done(); - + NotifierT * notifier) { + if (!get_flat_id()) { + notifier->store(NOTIFIER_OFFSET); + notifier->fence(); + } + notifier->sync(); + uint64_t offset_u64 {notifier->load()}; uint64_t raw_memory_u64 {reinterpret_cast(raw_memory)}; uint64_t address_u64 {raw_memory_u64 + offset_u64}; uint8_t* address {reinterpret_cast(address_u64)}; write_to_memory(address); - __syncthreads(); } -class NotifierTestFixture : public ::testing::Test { - using NotifierProxyT = NotifierProxy; - +class NotifierBase : public ::testing::Test { public: - NotifierTestFixture() { + NotifierBase() { assert(raw_memory_ == nullptr); hip_allocator_.allocate((void**)&raw_memory_, GIBIBYTE_); assert(raw_memory_); } - ~NotifierTestFixture() { + ~NotifierBase() { if (raw_memory_) { hip_allocator_.deallocate(raw_memory_); } } void - run_all_threads_once(uint32_t x_block_dim, - uint32_t x_grid_dim) { - const dim3 hip_blocksize(x_block_dim, 1, 1); - const dim3 hip_gridsize(x_grid_dim, 1, 1); - - hipLaunchKernelGGL(all_threads_once, - hip_gridsize, - hip_blocksize, - 0, - nullptr, - raw_memory_, - notifier_.get()); - - hipError_t return_code = hipStreamSynchronize(nullptr); - if (return_code != hipSuccess) { - printf("Failed in stream synchronize\n"); - assert(return_code == hipSuccess); - } - - size_t number_threads {x_block_dim * x_grid_dim}; - + verify(size_t number_threads) { uint8_t* offset_addr {compute_offset_addr()}; - for (size_t i {0}; i < number_threads; i++) { ASSERT_EQ(offset_addr[i], THREAD_VALUE); } @@ -136,12 +115,51 @@ class NotifierTestFixture : public ::testing::Test { */ uint8_t *raw_memory_ {nullptr}; +}; + +class NotifierBlockTestFixture : public NotifierBase { + using NotifierT = Notifier; + using NotifierProxyT = NotifierProxy; + + public: + void + run_all_threads_once(uint32_t x_block_dim, + uint32_t x_grid_dim) { + new (notifier_.get()) NotifierT(); + const dim3 block(x_block_dim, 1, 1); + const dim3 grid(x_grid_dim, 1, 1); + all_threads_once<<>>(raw_memory_, notifier_.get()); + CHECK_HIP(hipStreamSynchronize(nullptr)); + verify(x_block_dim * x_grid_dim); + } + /** * @brief Used to broadcast base offset for writing. */ NotifierProxyT notifier_ {}; }; +class NotifierAgentTestFixture : public NotifierBase { + using NotifierT = Notifier; + using NotifierProxyT = NotifierProxy; + + public: + void + run_all_threads_once(uint32_t x_block_dim, + uint32_t x_grid_dim) { + new (notifier_.get()) NotifierT(); + const dim3 block(x_block_dim, 1, 1); + const dim3 grid(x_grid_dim, 1, 1); + all_threads_once<<>>(raw_memory_, notifier_.get()); + CHECK_HIP(hipStreamSynchronize(nullptr)); + verify(x_block_dim * x_grid_dim); + } + + /** + * @brief Used to broadcast base offset for writing. + */ + NotifierProxyT notifier_ {}; +}; } // namespace rocshmem diff --git a/projects/rocshmem/tests/unit_tests/shmem_gtest.cpp b/projects/rocshmem/tests/unit_tests/shmem_gtest.cpp index 3115ef5e52..8536ca638a 100644 --- a/projects/rocshmem/tests/unit_tests/shmem_gtest.cpp +++ b/projects/rocshmem/tests/unit_tests/shmem_gtest.cpp @@ -21,8 +21,27 @@ *****************************************************************************/ #include "gtest/gtest.h" +#include int main(int argc, char **argv) { ::testing::InitGoogleTest(&argc, argv); - return RUN_ALL_TESTS(); + + int initialized; + MPI_Initialized(&initialized); + if (!initialized) { + int provided; + MPI_Init_thread(nullptr, nullptr, MPI_THREAD_MULTIPLE, &provided); + if (provided != MPI_THREAD_MULTIPLE) { + std::cerr << "MPI_THREAD_MULTIPLE support disabled.\n"; + } + } + + int ret_val = RUN_ALL_TESTS(); + + int finalized{0}; + MPI_Finalized(&finalized); + if (!finalized) { + MPI_Finalize(); + } + return ret_val; }