diff --git a/scripts/build_configs/ro_net_debug b/scripts/build_configs/ro_net_debug index 67c3f2d0a5..e8309b06d8 100755 --- a/scripts/build_configs/ro_net_debug +++ b/scripts/build_configs/ro_net_debug @@ -18,7 +18,7 @@ cmake \ -DPROFILE=OFF \ -DUSE_GPU_IB=OFF \ -DUSE_DC=OFF \ - -DUSE_IPC=OFF \ + -DUSE_IPC=ON \ -DUSE_THREADS=ON \ -DUSE_WF_COAL=OFF \ -DUSE_COHERENT_HEAP=ON \ diff --git a/src/util.hpp b/src/util.hpp index c209750d49..c02f891dc4 100644 --- a/src/util.hpp +++ b/src/util.hpp @@ -93,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. */ @@ -100,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. diff --git a/tests/unit_tests/ipc_impl_simple_coarse_gtest.hpp b/tests/unit_tests/ipc_impl_simple_coarse_gtest.hpp index 02dfd8c55a..083d73f418 100644 --- a/tests/unit_tests/ipc_impl_simple_coarse_gtest.hpp +++ b/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/tests/unit_tests/ipc_impl_simple_fine_gtest.hpp b/tests/unit_tests/ipc_impl_simple_fine_gtest.hpp index 3f85528e7d..fb60acdec5 100644 --- a/tests/unit_tests/ipc_impl_simple_fine_gtest.hpp +++ b/tests/unit_tests/ipc_impl_simple_fine_gtest.hpp @@ -33,9 +33,27 @@ namespace rocshmem { +enum TestType { + READ = 0, + WRITE = 1 +}; + __global__ void -kernel_simple_fine_copy(IpcImpl *ipc_impl, int *src, int *dest, size_t bytes) { +kernel_simple_fine_copy(IpcImpl *ipc_impl, int *src, int *dest, size_t bytes, TestType test) { + if (!threadIdx.x) { + ipc_impl->ipcCopy(dest, src, bytes); + ipc_impl->ipcFence(); + if (test == WRITE) { + ipc_impl->ipc + } + } + __syncthreads(); +} + +__global__ +void +kernel_simple_fine_copy_signal_validate(IpcImpl *ipc_impl, int *src, int *dest, size_t bytes) { if (!threadIdx.x) { ipc_impl->ipcCopy(dest, src, bytes); ipc_impl->ipcFence(); @@ -51,6 +69,14 @@ kernel_simple_fine_copy_wg(IpcImpl *ipc_impl, int *src, int *dest, size_t bytes) __syncthreads(); } +__global__ +void +kernel_simple_fine_copy_wg_signal_validate(IpcImpl *ipc_impl, int *src, int *dest, size_t bytes) { + ipc_impl->ipcCopy_wg(dest, src, bytes); + ipc_impl->ipcFence(); + __syncthreads(); +} + __global__ void kernel_simple_fine_copy_wave(IpcImpl *ipc_impl, int *src, int *dest, size_t bytes) { @@ -59,6 +85,14 @@ kernel_simple_fine_copy_wave(IpcImpl *ipc_impl, int *src, int *dest, size_t byte __syncthreads(); } +__global__ +void +kernel_simple_fine_copy_wave_signal_validate(IpcImpl *ipc_impl, int *src, int *dest, size_t bytes) { + ipc_impl->ipcCopy_wave(dest, src, bytes); + ipc_impl->ipcFence(); + __syncthreads(); +} + class IPCImplSimpleFineTestFixture : public ::testing::Test { using HEAP_T = HeapMemory; @@ -91,51 +125,46 @@ class IPCImplSimpleFineTestFixture : public ::testing::Test { CHECK_HIP(hipStreamSynchronize(nullptr)); } - enum TestType { - READ = 0, - WRITE = 1 - }; - void write(const dim3 grid, const dim3 block, size_t elems) { iota_golden(elems); initialize_src_buffer(WRITE); copy(WRITE, grid, block); - validate_dest_buffer(WRITE); + check_device_validation_errors(WRITE); } void write_wg(const dim3 grid, const dim3 block, size_t elems) { iota_golden(elems); initialize_src_buffer(WRITE); copy_wg(WRITE, grid, block); - validate_dest_buffer(WRITE); + check_device_validation_errors(WRITE); } void write_wave(const dim3 grid, const dim3 block, size_t elems) { iota_golden(elems); initialize_src_buffer(WRITE); copy_wave(WRITE, grid, block); - validate_dest_buffer(WRITE); + check_device_validation_errors(WRITE); } void read(const dim3 grid, const dim3 block, size_t elems) { iota_golden(elems); initialize_src_buffer(READ); copy(READ, grid, block); - validate_dest_buffer(READ); + check_device_validation_errors(READ); } void read_wg(const dim3 grid, const dim3 block, size_t elems) { iota_golden(elems); initialize_src_buffer(READ); copy_wg(READ, grid, block); - validate_dest_buffer(READ); + check_device_validation_errors(READ); } void read_wave(const dim3 grid, const dim3 block, size_t elems) { iota_golden(elems); initialize_src_buffer(READ); copy_wave(READ, grid, block); - validate_dest_buffer(READ); + check_device_validation_errors(READ); } void iota_golden(size_t elems) { @@ -160,6 +189,7 @@ class IPCImplSimpleFineTestFixture : public ::testing::Test { CHECK_HIP(hipStreamSynchronize(nullptr)); } + __host__ __device__ bool pe_initializes_src_buffer(TestType test) { bool is_write_test = test; bool is_read_test = !test; @@ -184,7 +214,7 @@ class IPCImplSimpleFineTestFixture : public ::testing::Test { } size_t bytes = golden_.size() * sizeof(int); mpi_.barrier(); - launch(fn, grid, block, src, dest, bytes); + launch(fn, grid, block, src, dest, bytes, test); mpi_.barrier(); } @@ -200,6 +230,13 @@ class IPCImplSimpleFineTestFixture : public ::testing::Test { 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(validation_error, false); + } + void validate_dest_buffer(TestType test) { if (!pe_validates_dest_buffer(test)) { return; @@ -211,6 +248,21 @@ class IPCImplSimpleFineTestFixture : public ::testing::Test { } } + __device__ + 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 {get_flat_id()}; i < golden_.size(); i += get_flat_grid_size()) { + if (dev_golden_[i] != dev_dest[i]) { + validation_error = true; + } + } + } + + __host__ __device__ bool pe_validates_dest_buffer(TestType test) { return !pe_initializes_src_buffer(test); } @@ -218,7 +270,7 @@ class IPCImplSimpleFineTestFixture : public ::testing::Test { protected: std::vector golden_; - std::vector output_; + std::vector device_golden_; HEAP_T heap_mem_ {}; @@ -229,6 +281,8 @@ class IPCImplSimpleFineTestFixture : public ::testing::Test { IpcImpl *ipc_impl_dptr_ {nullptr}; HIPAllocator hip_allocator_ {}; + + bool validation_error {false}; }; } // namespace rocshmem