diff --git a/projects/rocshmem/tests/unit_tests/CMakeLists.txt b/projects/rocshmem/tests/unit_tests/CMakeLists.txt index b47da1c154..e3ce4498ad 100644 --- a/projects/rocshmem/tests/unit_tests/CMakeLists.txt +++ b/projects/rocshmem/tests/unit_tests/CMakeLists.txt @@ -91,6 +91,7 @@ target_sources( #context_ipc_gtest.cpp ipc_impl_simple_coarse_gtest.cpp ipc_impl_simple_fine_gtest.cpp + ipc_impl_tiled_fine_gtest.cpp ) ############################################################################### 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 index 6d7b758e31..e2492fafad 100644 --- a/projects/rocshmem/tests/unit_tests/ipc_impl_simple_fine_gtest.hpp +++ b/projects/rocshmem/tests/unit_tests/ipc_impl_simple_fine_gtest.hpp @@ -45,7 +45,7 @@ enum TestType { __device__ void -validator(bool *error, int *golden, int *dest, size_t bytes) { +simple_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]) { @@ -58,7 +58,7 @@ validator(bool *error, int *golden, int *dest, size_t bytes) { template __global__ void -kernel_put_with_signal_validator(bool *error, int *golden, int *dest, size_t bytes, NotifierT *notifier) { +kernel_put_with_signal_simple_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) { @@ -66,7 +66,7 @@ kernel_put_with_signal_validator(bool *error, int *golden, int *dest, size_t byt } } notifier->sync(); - validator(error, golden, dest, bytes); + simple_validator(error, golden, dest, bytes); } template @@ -82,7 +82,7 @@ kernel_simple_fine_copy(IpcImpl *ipc_impl, bool *error, int *golden, int *src, i } if (test == READ) { notifier->sync(); - validator(error, golden, dest, bytes); + simple_validator(error, golden, dest, bytes); } } @@ -101,7 +101,7 @@ kernel_simple_fine_copy_block(IpcImpl *ipc_impl, bool *error, int *golden, int * } if (test == READ) { notifier->sync(); - validator(error, golden, dest, bytes); + simple_validator(error, golden, dest, bytes); } } @@ -121,7 +121,7 @@ kernel_simple_fine_copy_warp(IpcImpl *ipc_impl, bool *error, int *golden, int *s __syncthreads(); if (test == READ) { notifier->sync(); - validator(error, golden, dest, bytes); + simple_validator(error, golden, dest, bytes); } } @@ -237,7 +237,7 @@ class IPCImplSimpleFine : public ::testing::TestWithParam(ipc_impl_.ipc_bases[1]); - FN_T2 val_fn = kernel_put_with_signal_validator; + FN_T2 val_fn = kernel_put_with_signal_simple_validator; launch(val_fn, grid, block, dest, bytes); } mpi_.barrier(); diff --git a/projects/rocshmem/tests/unit_tests/ipc_impl_tiled_fine_gtest.cpp b/projects/rocshmem/tests/unit_tests/ipc_impl_tiled_fine_gtest.cpp new file mode 100644 index 0000000000..9e4ecb6846 --- /dev/null +++ b/projects/rocshmem/tests/unit_tests/ipc_impl_tiled_fine_gtest.cpp @@ -0,0 +1,321 @@ +/****************************************************************************** + * 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_tiled_fine_gtest.hpp" + +using namespace rocshmem; + +//============================================================================= + +TEST_P(DegenerateTiledFine, ptr_check) { + ASSERT_NE(heap_mem_.get_ptr(), nullptr); +} + +TEST_P(DegenerateTiledFine, MPI_num_pes) { + ASSERT_EQ(mpi_.num_pes(), 2); +} + +TEST_P(DegenerateTiledFine, IPC_bases) { + ASSERT_EQ(mpi_.num_pes(), 2); + 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_P(DegenerateTiledFine, golden_1048576_int) { + iota_golden(1048576); + validate_golden(1048576); +} + +INSTANTIATE_TEST_SUITE_P( + IPCImplTiledFineTestFixture, + DegenerateTiledFine, + ::testing::Values( + std::make_tuple(1, 1, 1)) +); + +//============================================================================= + +int block_signals_calculation(int grid_dim_x, int block_dim_x, size_t size) { + size_t bytes = size * sizeof(int); + int total_num_threads {grid_dim_x * block_dim_x}; + int one_grid_iteration_data_size {total_num_threads * THREAD_TRANSFER_GRANULARITY}; + int num_grid_iterations = ((bytes - THREAD_TRANSFER_GRANULARITY) + one_grid_iteration_data_size) / one_grid_iteration_data_size; + int partial_grid_last_iteration_data_size = bytes % one_grid_iteration_data_size; + + int bytes_per_block {block_dim_x * THREAD_TRANSFER_GRANULARITY}; + + int num_signals_for_one_full_iteration = grid_dim_x; + int num_signals_for_partial_last_iteration = ((partial_grid_last_iteration_data_size - THREAD_TRANSFER_GRANULARITY) + bytes_per_block) / bytes_per_block; + + int num_signals = 0; + if (partial_grid_last_iteration_data_size) { + num_signals = num_signals_for_one_full_iteration * (num_grid_iterations - 1); + num_signals += num_signals_for_partial_last_iteration; + } else { + num_signals = num_signals_for_one_full_iteration * num_grid_iterations; + } + + return num_signals; +} + +TEST_P(ParameterizedBlockTiledFine, write) { + dim3 grid = dim3(std::get<0>(GetParam()), 1, 1); + dim3 block = dim3(std::get<1>(GetParam()), 1, 1); + size_t size = std::get<2>(GetParam()); + auto number_of_signals_required = block_signals_calculation(std::get<0>(GetParam()), std::get<1>(GetParam()), size); + write(grid, block, size, number_of_signals_required); +} + +TEST_P(ParameterizedBlockTiledFine, read) { + dim3 grid = dim3(std::get<0>(GetParam()), 1, 1); + dim3 block = dim3(std::get<1>(GetParam()), 1, 1); + size_t size = std::get<2>(GetParam()); + read(grid, block, size); +} + +INSTANTIATE_TEST_SUITE_P( + IPCImplTiledFineTestFixture, + ParameterizedBlockTiledFine, + ::testing::Values( + std::make_tuple(1, 1024, 32), // 0 + std::make_tuple(1, 1024, 2048), // 1 + std::make_tuple(1, 1024, 65536), // 2 + std::make_tuple(1, 1, 1048576), // 3 + std::make_tuple(1, 2, 1048576), // 4 + std::make_tuple(1, 4, 1048576), // 5 + std::make_tuple(1, 8, 1048576), // 6 + std::make_tuple(1, 16, 1048576), // 7 + std::make_tuple(1, 32, 1048576), // 8 + std::make_tuple(1, 64, 1048576), // 9 + std::make_tuple(1, 128, 1048576), // 10 + std::make_tuple(1, 256, 1048576), // 11 + std::make_tuple(1, 512, 1048576), // 12 + std::make_tuple(1, 768, 1048576), // 13 + std::make_tuple(1, 1024, 1048576), // 14 + std::make_tuple(2, 1, 1048576), // 15 + std::make_tuple(2, 2, 1048576), // 16 + std::make_tuple(2, 4, 1048576), // 17 + std::make_tuple(2, 8, 1048576), // 18 + std::make_tuple(2, 16, 1048576), // 19 + std::make_tuple(2, 32, 1048576), // 20 + std::make_tuple(2, 64, 1048576), // 21 + std::make_tuple(2, 128, 1048576), // 22 + std::make_tuple(2, 256, 1048576), // 23 + std::make_tuple(2, 512, 1048576), // 24 + std::make_tuple(2, 768, 1048576), // 25 + std::make_tuple(2, 1024, 1048576), // 26 + std::make_tuple(4, 1024, 1048576), // 27 + std::make_tuple(8, 1024, 1048576), // 28 + std::make_tuple(16, 1024, 1048576), // 29 + std::make_tuple(32, 1024, 1048576), // 30 + std::make_tuple(38, 1024, 1048576), // 31 + std::make_tuple(38, 1024, 2097152), // 32 + std::make_tuple(38, 1024, 4194304), // 33 + std::make_tuple(38, 1024, 8388608), // 34 + std::make_tuple(38, 1024, 16777216), // 35 + std::make_tuple(38, 1024, 33554432)) // 36 +); + +//============================================================================= + +int warp_signals_calculation(int grid_dim_x, int block_dim_x, size_t size) { + size_t bytes = size * sizeof(int); + int total_num_threads {grid_dim_x * block_dim_x}; + int one_grid_iteration_data_size {total_num_threads * THREAD_TRANSFER_GRANULARITY}; + int num_grid_iterations = ((bytes - THREAD_TRANSFER_GRANULARITY) + one_grid_iteration_data_size) / one_grid_iteration_data_size; + int partial_grid_last_iteration_data_size = bytes % one_grid_iteration_data_size; + + int warps_per_block {warpsPerBlock(block_dim_x)}; + int total_num_warps_in_grid {grid_dim_x * warps_per_block}; + int bytes_per_warp {WARP_SIZE * THREAD_TRANSFER_GRANULARITY}; + + int num_signals_for_one_full_iteration = total_num_warps_in_grid; + int num_signals_for_partial_last_iteration = ((partial_grid_last_iteration_data_size - THREAD_TRANSFER_GRANULARITY) + bytes_per_warp) / bytes_per_warp; + + int num_signals = 0; + if (partial_grid_last_iteration_data_size) { + num_signals = num_signals_for_one_full_iteration * (num_grid_iterations - 1); + num_signals += num_signals_for_partial_last_iteration; + } else { + num_signals = num_signals_for_one_full_iteration * num_grid_iterations; + } + + return num_signals; +} + +TEST_P(ParameterizedWarpTiledFine, write) { + dim3 grid = dim3(std::get<0>(GetParam()), 1, 1); + dim3 block = dim3(std::get<1>(GetParam()), 1, 1); + size_t size = std::get<2>(GetParam()); + auto number_of_signals_required = warp_signals_calculation(std::get<0>(GetParam()), std::get<1>(GetParam()), size); + write(grid, block, size, number_of_signals_required); +} + +TEST_P(ParameterizedWarpTiledFine, read) { + dim3 grid = dim3(std::get<0>(GetParam()), 1, 1); + dim3 block = dim3(std::get<1>(GetParam()), 1, 1); + size_t size = std::get<2>(GetParam()); + read(grid, block, size); +} + +INSTANTIATE_TEST_SUITE_P( + IPCImplTiledFineTestFixture, + ParameterizedWarpTiledFine, + ::testing::Values( + std::make_tuple(1, 64, 8), // 0 + std::make_tuple(1, 64, 32), // 1 + std::make_tuple(1, 1, 1048576), // 2 + std::make_tuple(1, 2, 1048576), // 3 + std::make_tuple(1, 3, 1048576), // 4 + std::make_tuple(1, 4, 1048576), // 5 + std::make_tuple(1, 5, 1048576), // 6 + std::make_tuple(1, 6, 1048576), // 7 + std::make_tuple(1, 7, 1048576), // 8 + std::make_tuple(1, 8, 1048576), // 9 + std::make_tuple(1, 9, 1048576), // 10 + std::make_tuple(1, 10, 1048576), // 11 + std::make_tuple(1, 11, 1048576), // 12 + std::make_tuple(1, 12, 1048576), // 13 + std::make_tuple(1, 13, 1048576), // 14 + std::make_tuple(1, 14, 1048576), // 15 + std::make_tuple(1, 15, 1048576), // 16 + std::make_tuple(1, 16, 1048576), // 17 + std::make_tuple(1, 17, 1048576), // 18 + std::make_tuple(1, 18, 1048576), // 19 + std::make_tuple(1, 19, 1048576), // 20 + std::make_tuple(1, 20, 1048576), // 21 + std::make_tuple(1, 21, 1048576), // 22 + std::make_tuple(1, 22, 1048576), // 23 + std::make_tuple(1, 23, 1048576), // 24 + std::make_tuple(1, 24, 1048576), // 25 + std::make_tuple(1, 25, 1048576), // 26 + std::make_tuple(1, 26, 1048576), // 27 + std::make_tuple(1, 27, 1048576), // 28 + std::make_tuple(1, 28, 1048576), // 29 + std::make_tuple(1, 28, 1048576), // 30 + std::make_tuple(1, 29, 1048576), // 31 + std::make_tuple(1, 30, 1048576), // 32 + std::make_tuple(1, 31, 1048576), // 33 + std::make_tuple(1, 32, 1048576), // 34 + std::make_tuple(1, 33, 1048576), // 35 + std::make_tuple(1, 34, 1048576), // 36 + std::make_tuple(1, 35, 1048576), // 37 + std::make_tuple(1, 36, 1048576), // 38 + std::make_tuple(1, 37, 1048576), // 39 + std::make_tuple(1, 38, 1048576), // 40 + std::make_tuple(1, 39, 1048576), // 41 + std::make_tuple(1, 40, 1048576), // 42 + std::make_tuple(1, 41, 1048576), // 43 + std::make_tuple(1, 42, 1048576), // 44 + std::make_tuple(1, 43, 1048576), // 45 + std::make_tuple(1, 44, 1048576), // 46 + std::make_tuple(1, 45, 1048576), // 47 + std::make_tuple(1, 46, 1048576), // 48 + std::make_tuple(1, 47, 1048576), // 49 + std::make_tuple(1, 48, 1048576), // 50 + std::make_tuple(1, 49, 1048576), // 51 + std::make_tuple(1, 50, 1048576), // 52 + std::make_tuple(1, 51, 1048576), // 53 + std::make_tuple(1, 52, 1048576), // 54 + std::make_tuple(1, 53, 1048576), // 55 + std::make_tuple(1, 54, 1048576), // 56 + std::make_tuple(1, 55, 1048576), // 57 + std::make_tuple(1, 56, 1048576), // 58 + std::make_tuple(1, 57, 1048576), // 59 + std::make_tuple(1, 58, 1048576), // 60 + std::make_tuple(1, 59, 1048576), // 61 + std::make_tuple(1, 60, 1048576), // 62 + std::make_tuple(1, 61, 1048576), // 63 + std::make_tuple(1, 62, 1048576), // 64 + std::make_tuple(1, 63, 1048576), // 65 + std::make_tuple(1, 64, 1048576), // 66 + std::make_tuple(1, 256, 1048576), // 67 + std::make_tuple(1, 512, 1048576), // 68 + std::make_tuple(1, 768, 1048576), // 69 + std::make_tuple(1, 1024, 1048576), // 70 + std::make_tuple(2, 32, 1048576), // 71 + std::make_tuple(2, 64, 1048576), // 72 + std::make_tuple(2, 128, 1048576), // 73 + std::make_tuple(2, 256, 1048576), // 74 + std::make_tuple(2, 512, 1048576), // 75 + std::make_tuple(2, 1024, 1048576), // 76 + std::make_tuple(4, 32, 1048576), // 77 + std::make_tuple(4, 64, 1048576), // 78 + std::make_tuple(4, 128, 1048576), // 79 + std::make_tuple(4, 256, 1048576), // 80 + std::make_tuple(4, 512, 1048576), // 81 + std::make_tuple(4, 1024, 1048576), // 82 + std::make_tuple(8, 32, 1048576), // 83 + std::make_tuple(8, 64, 1048576), // 84 + std::make_tuple(8, 128, 1048576), // 85 + std::make_tuple(8, 256, 1048576), // 86 + std::make_tuple(8, 512, 1048576), // 87 + std::make_tuple(8, 1024, 1048576), // 88 + std::make_tuple(16, 32, 1048576), // 89 + std::make_tuple(16, 64, 1048576), // 90 + std::make_tuple(16, 128, 1048576), // 91 + std::make_tuple(16, 256, 1048576), // 92 + std::make_tuple(16, 512, 1048576), // 93 + std::make_tuple(16, 1024, 1048576), // 94 + std::make_tuple(32, 32, 1048576), // 95 + std::make_tuple(32, 64, 1048576), // 96 + std::make_tuple(32, 128, 1048576), // 97 + std::make_tuple(32, 256, 1048576), // 98 + std::make_tuple(32, 512, 1048576), // 99 + std::make_tuple(32, 1024, 1048576), // 100 + std::make_tuple(38, 32, 1048576), // 101 + std::make_tuple(38, 64, 1048576), // 102 + std::make_tuple(38, 128, 1048576), // 103 + std::make_tuple(38, 256, 1048576), // 104 + std::make_tuple(38, 512, 1048576), // 105 + std::make_tuple(38, 1024, 1048576), // 106 + std::make_tuple(38, 1024, 2097152), // 107 + std::make_tuple(38, 1024, 4194304), // 108 + std::make_tuple(38, 1024, 8388608), // 109 + std::make_tuple(38, 1024, 16777216), // 110 + std::make_tuple(38, 1024, 33554432)) // 111 +); + +//============================================================================= + +TEST_P(ParameterizedThreadTiledFine, write) { + dim3 grid = dim3(std::get<0>(GetParam()), 1, 1); + dim3 block = dim3(std::get<1>(GetParam()), 1, 1); + size_t size = std::get<2>(GetParam()); + write(grid, block, size, 1); +} + +TEST_P(ParameterizedThreadTiledFine, read) { + dim3 grid = dim3(std::get<0>(GetParam()), 1, 1); + dim3 block = dim3(std::get<1>(GetParam()), 1, 1); + size_t size = std::get<2>(GetParam()); + read(grid, block, size); +} + +INSTANTIATE_TEST_SUITE_P( + IPCImplTiledFineTestFixture, + ParameterizedThreadTiledFine, + ::testing::Values( + std::make_tuple(1, 1, 1048576)) +); diff --git a/projects/rocshmem/tests/unit_tests/ipc_impl_tiled_fine_gtest.hpp b/projects/rocshmem/tests/unit_tests/ipc_impl_tiled_fine_gtest.hpp new file mode 100644 index 0000000000..ea257b11e2 --- /dev/null +++ b/projects/rocshmem/tests/unit_tests/ipc_impl_tiled_fine_gtest.hpp @@ -0,0 +1,354 @@ +/****************************************************************************** + * 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 { + +const int WARP_SIZE = 64; + +const int THREAD_TRANSFER_GRANULARITY = 8; // DWORDX2 + +int warpsPerBlock(size_t block_size) { + return ((block_size - 1) + WARP_SIZE) / WARP_SIZE; +} + +// set signal pointer to ipc_impl.ipc_bases location at unused offset +// be careful with test size to not overrun this location +const uint32_t SIGNAL_OFFSET {67108864}; + +enum TestType { + READ = 0, + WRITE = 1 +}; + +__device__ +void +tiled_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_tiled_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(); + tiled_validator(error, golden, dest, bytes); +} + +template +__global__ +void +kernel_tiled_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(); + tiled_validator(error, golden, dest, bytes); + } +} + +template +__global__ +void +kernel_tiled_fine_copy_block(IpcImpl *ipc_impl, bool *error, int *golden, int *src, int *dest, size_t bytes, TestType test, NotifierT *notifier) { + int block_bytes = blockDim.x * THREAD_TRANSFER_GRANULARITY; + int block_byte_offset = blockIdx.x * block_bytes; + for (int i {block_byte_offset}; i < bytes; i += get_flat_grid_size() * THREAD_TRANSFER_GRANULARITY) { + int chunk = min(block_bytes, bytes - i); + ipc_impl->ipcCopy_wg((char*)dest + i, (char*)src + i, chunk); + ipc_impl->ipcFence(); + __syncthreads(); + if (test == WRITE) { + if (!threadIdx.x) { + ipc_impl->ipcAMOFetchAdd(dest + SIGNAL_OFFSET, -1); + } + } + } + if (test == READ) { + notifier->sync(); + tiled_validator(error, golden, dest, bytes); + } +} + +template +__global__ +void +kernel_tiled_fine_copy_warp(IpcImpl *ipc_impl, bool *error, int *golden, int *src, int *dest, size_t bytes, TestType test, NotifierT *notifier) { + int warp_id = (blockIdx.x * blockDim.x + threadIdx.x) / WARP_SIZE; + int warp_bytes = WARP_SIZE * THREAD_TRANSFER_GRANULARITY; + int warp_byte_offset = warp_id * warp_bytes; + for (int i {warp_byte_offset}; i < bytes; i += get_flat_grid_size() * THREAD_TRANSFER_GRANULARITY) { + int chunk = min(warp_bytes, bytes - i); + ipc_impl->ipcCopy_wave(((char*)dest) + i, ((char*)src) + i, chunk); + ipc_impl->ipcFence(); + if (test == WRITE) { + if (!(threadIdx.x % WARP_SIZE)) { + ipc_impl->ipcAMOFetchAdd(dest + SIGNAL_OFFSET, -1); + } + } + } + __syncthreads(); + if (test == READ) { + notifier->sync(); + tiled_validator(error, golden, dest, bytes); + } +} + +class IPCImplTiledFine : public ::testing::TestWithParam> { + 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: + IPCImplTiledFine() { + 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; + } + + ~IPCImplTiledFine() { + 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)); + } + + virtual void copy(TestType test, dim3 grid, dim3 block) { + FAIL(); + } + + void write(const dim3 grid, const dim3 block, size_t elems, int signal_value) { + iota_golden(elems); + initialize_signal(WRITE, signal_value); + initialize_src_buffer(WRITE); + copy(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 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, int signal_value = 0) { + 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) = signal_value; + } + } + + 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_tiled_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 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}; +}; + +class DegenerateTiledFine : public IPCImplTiledFine { + public: + ~DegenerateTiledFine() override {}; +}; + +class ParameterizedBlockTiledFine : public IPCImplTiledFine { + public: + ~ParameterizedBlockTiledFine() override {}; + + void copy(TestType test, dim3 grid, dim3 block) override { + execute(test, kernel_tiled_fine_copy_block, grid, block); + } +}; + +class ParameterizedWarpTiledFine : public IPCImplTiledFine { + public: + ~ParameterizedWarpTiledFine() override {}; + + void copy(TestType test, dim3 grid, dim3 block) override { + execute(test, kernel_tiled_fine_copy_warp, grid, block); + } +}; + +class ParameterizedThreadTiledFine : public IPCImplTiledFine { + public: + ~ParameterizedThreadTiledFine() override {}; + + void copy(TestType test, dim3 grid, dim3 block) override { + execute(test, kernel_tiled_fine_copy, grid, block); + } +}; + +} // namespace rocshmem + +#endif // ROCSHMEM_IPC_IMPL_SIMPLE_FINE_GTEST_HPP