From 68384da0199c3ff5bbe87fdb9210a5b8b67df3c5 Mon Sep 17 00:00:00 2001 From: Brandon Potter Date: Tue, 13 Aug 2024 12:01:24 -0700 Subject: [PATCH] Update Notifier fixture to Block [ROCm/rocshmem commit: 5b42cff96cea239991c144cfbb84f1abd4d5c5bb] --- projects/rocshmem/src/atomic.hpp | 12 +++++++- projects/rocshmem/src/memory/notifier.hpp | 10 +++++-- .../tests/unit_tests/notifier_gtest.cpp | 28 +++++++++---------- .../tests/unit_tests/notifier_gtest.hpp | 23 ++++++++------- 4 files changed, 43 insertions(+), 30 deletions(-) diff --git a/projects/rocshmem/src/atomic.hpp b/projects/rocshmem/src/atomic.hpp index 5d9dc30886..330bd82f86 100644 --- a/projects/rocshmem/src/atomic.hpp +++ b/projects/rocshmem/src/atomic.hpp @@ -122,7 +122,7 @@ T fetch_min(T* obj, U arg, rocshmem_memory_orders o) { template __device__ -void thread_fence() { +void threadfence() { if constexpr (s == memory_scope_system) { __threadfence_system(); } else if constexpr (s == memory_scope_agent) { @@ -132,6 +132,16 @@ void thread_fence() { } } +template +__device__ +void sync() { + if constexpr (s == memory_scope_workgroup) { + __syncthreads(); + } else { + assert(false); + } +} + } // namespace atomic } // namespace detail } // namespace rocshmem diff --git a/projects/rocshmem/src/memory/notifier.hpp b/projects/rocshmem/src/memory/notifier.hpp index 1f9f6efe76..12f53a4c8b 100644 --- a/projects/rocshmem/src/memory/notifier.hpp +++ b/projects/rocshmem/src/memory/notifier.hpp @@ -33,15 +33,19 @@ template class Notifier { public: __device__ uint64_t load() { - return detail::atomic::load(&value_, orders); + return detail::atomic::load(&value_, orders); } __device__ void store(uint64_t val) { - detail::atomic::store(&value_, val, orders); + detail::atomic::store(&value_, val, orders); } __device__ void fence() { - detail::atomic::thread_fence(); + detail::atomic::threadfence(); + } + + __device__ void sync() { + detail::atomic::sync(); } private: diff --git a/projects/rocshmem/tests/unit_tests/notifier_gtest.cpp b/projects/rocshmem/tests/unit_tests/notifier_gtest.cpp index 9f79c62397..e1ebea0210 100644 --- a/projects/rocshmem/tests/unit_tests/notifier_gtest.cpp +++ b/projects/rocshmem/tests/unit_tests/notifier_gtest.cpp @@ -28,30 +28,30 @@ using namespace rocshmem; ******************************* Fixture Tests ******************************* *****************************************************************************/ -TEST_F(NotifierTestFixture, run_all_threads_once_1_1) { - run_all_threads_once(1, 1); +TEST_F(NotifierBlockTestFixture, run_all_threads_once_1_1) { + run_all_threads_once_block(1, 1); } -TEST_F(NotifierTestFixture, run_all_threads_once_2_1) { - run_all_threads_once(2, 1); +TEST_F(NotifierBlockTestFixture, run_all_threads_once_2_1) { + run_all_threads_once_block(2, 1); } -TEST_F(NotifierTestFixture, run_all_threads_once_64_1) { - run_all_threads_once(64, 1); +TEST_F(NotifierBlockTestFixture, run_all_threads_once_64_1) { + run_all_threads_once_block(64, 1); } -TEST_F(NotifierTestFixture, run_all_threads_once_128_1) { - run_all_threads_once(128, 1); +TEST_F(NotifierBlockTestFixture, run_all_threads_once_128_1) { + run_all_threads_once_block(128, 1); } -TEST_F(NotifierTestFixture, run_all_threads_once_256_1) { - run_all_threads_once(256, 1); +TEST_F(NotifierBlockTestFixture, run_all_threads_once_256_1) { + run_all_threads_once_block(256, 1); } -TEST_F(NotifierTestFixture, run_all_threads_once_512_1) { - run_all_threads_once(512, 1); +TEST_F(NotifierBlockTestFixture, run_all_threads_once_512_1) { + run_all_threads_once_block(512, 1); } -TEST_F(NotifierTestFixture, run_all_threads_once_1024_1) { - run_all_threads_once(1024, 1); +TEST_F(NotifierBlockTestFixture, run_all_threads_once_1024_1) { + run_all_threads_once_block(1024, 1); } diff --git a/projects/rocshmem/tests/unit_tests/notifier_gtest.hpp b/projects/rocshmem/tests/unit_tests/notifier_gtest.hpp index 303f1932c9..453b93680d 100644 --- a/projects/rocshmem/tests/unit_tests/notifier_gtest.hpp +++ b/projects/rocshmem/tests/unit_tests/notifier_gtest.hpp @@ -43,7 +43,7 @@ static const uint64_t NOTIFIER_OFFSET {0x100B00}; inline __device__ void -write_to_memory(uint8_t* raw_memory) { +write_to_memory_block(uint8_t* raw_memory) { auto thread_idx {get_flat_block_id()}; raw_memory[thread_idx] = THREAD_VALUE; __threadfence(); @@ -51,44 +51,43 @@ write_to_memory(uint8_t* raw_memory) { __global__ void -all_threads_once(uint8_t* raw_memory, - Notifier * notifier) { +all_threads_once_block(uint8_t* raw_memory, + Notifier * notifier) { if (!threadIdx.x) { notifier->store(NOTIFIER_OFFSET); notifier->fence(); } - __syncthreads(); + 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(); + write_to_memory_block(address); } -class NotifierTestFixture : public ::testing::Test { +class NotifierBlockTestFixture : public ::testing::Test { using NotifierProxyT = NotifierProxy; public: - NotifierTestFixture() { + NotifierBlockTestFixture() { assert(raw_memory_ == nullptr); hip_allocator_.allocate((void**)&raw_memory_, GIBIBYTE_); assert(raw_memory_); } - ~NotifierTestFixture() { + ~NotifierBlockTestFixture() { if (raw_memory_) { hip_allocator_.deallocate(raw_memory_); } } void - run_all_threads_once(uint32_t x_block_dim, - uint32_t x_grid_dim) { + run_all_threads_once_block(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, + hipLaunchKernelGGL(all_threads_once_block, hip_gridsize, hip_blocksize, 0,