Add tiled fine-grained unit tests

[ROCm/rocshmem commit: d241015e0f]
2024-11-04 17:16:07 -06:00
@@ -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
 )

 ###############################################################################
@@ -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 <typename NotifierT>
 __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<int, detail::atomic::memory_scope_system>(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 <typename NotifierT>
@@ -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<std::tuple<int, int, i
            mpi_.barrier();
            if (test == WRITE) {
                int *dest = reinterpret_cast<int*>(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();
@@ -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))
+);
@@ -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 <numeric>
+#include <mpi.h>
+
+#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 <typename NotifierT>
+__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<int, detail::atomic::memory_scope_system>(dest + SIGNAL_OFFSET, orders) != 0) {
+            ;
+        }
+    }
+    notifier->sync();
+    tiled_validator(error, golden, dest, bytes);
+}
+
+template <typename NotifierT>
+__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 <typename NotifierT>
+__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 <typename NotifierT>
+__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<std::tuple<int, int, int>> {
+    using HEAP_T = HeapMemory<HIPDefaultFinegrainedAllocator>;
+    using MPI_T = RemoteHeapInfo<CommunicatorMPI>;
+    using NotifierT = Notifier<detail::atomic::memory_scope_agent>;
+    using NotifierProxyT = NotifierProxy<HIPAllocator, detail::atomic::memory_scope_agent>;
+    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<<<grid, block>>>(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<<<grid, block>>>(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<int*>(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<int*>(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<int*>(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<int*>(ipc_impl_.ipc_bases[0]);
+            dest = reinterpret_cast<int*>(ipc_impl_.ipc_bases[1]);
+        } else {
+            src = reinterpret_cast<int*>(ipc_impl_.ipc_bases[1]);
+            dest = reinterpret_cast<int*>(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<int*>(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<int> 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