Add AtomicWFQueue implementation and tests (#62)

* feat: Add AtomicWFQueue implementation - Implemented wavefront-safe atomic FIFO queue ensuring first-come, first-serve order - Added efficient synchronization using atomics - Enhanced `dequeue` to wait until an element is available * test: Add GTest for AtomicWFQueue - Implemented unit tests for AtomicWFQueue using GoogleTest framework - Added tests for `enqueue`, `dequeue`, and edge cases - Ensured synchronization behavior and correctness under concurrent conditions * Add assert in `enqueue` and update atomics - Added an assert in the `enqueue` function to ensure it fails if the queue is full
2025-03-25 00:45:19 -05:00
parent c16b0d6952
commit b84b5638cf
5 changed files with 891 additions and 0 deletions
@@ -92,6 +92,7 @@ target_sources(
    free_list_gtest.cpp
    #context_ipc_gtest.cpp
    wavefront_size_gtest.cpp
+    atomic_wf_queue_gtest.cpp
 )

 if (USE_IPC)
@@ -0,0 +1,199 @@
+/******************************************************************************
+ * 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 "atomic_wf_queue_gtest.hpp"
+
+using namespace rocshmem;
+
+/*****************************************************************************
+ ******************************* Fixture Tests *******************************
+ *****************************************************************************/
+
+TEST_F(AtomicWFQueueTestFixture, active_logical_lane_ids_4) {
+  get_thread_lane_ids(4);
+}
+
+TEST_F(AtomicWFQueueTestFixture, active_logical_lane_ids_10) {
+  get_thread_lane_ids(10);
+}
+
+TEST_F(AtomicWFQueueTestFixture, active_logical_lane_ids_16) {
+  get_thread_lane_ids(16);
+}
+
+TEST_F(AtomicWFQueueTestFixture, active_logical_lane_ids_17) {
+  get_thread_lane_ids(17);
+}
+
+TEST_F(AtomicWFQueueTestFixture, active_logical_lane_ids_64) {
+  get_thread_lane_ids(64);
+}
+
+TEST_F(AtomicWFQueueTestFixture, active_logical_lane_ids_97) {
+  get_thread_lane_ids(97);
+}
+
+TEST_F(AtomicWFQueueTestFixture, active_logical_lane_ids_256) {
+  get_thread_lane_ids(256);
+}
+
+TEST_F(AtomicWFQueueTestFixture, active_logical_lane_ids_183) {
+  get_thread_lane_ids(183);
+}
+
+TEST_F(AtomicWFQueueTestFixture, active_logical_lane_ids_1024) {
+  get_thread_lane_ids(1024);
+}
+
+TEST_F(AtomicWFQueueTestFixture, init_2_64) {
+  init_queue(2, 64);
+}
+
+TEST_F(AtomicWFQueueTestFixture, init_64_256) {
+  init_queue(64, 256);
+}
+
+TEST_F(AtomicWFQueueTestFixture, init_1024_96) {
+  init_queue(1024, 96);
+}
+
+TEST_F(AtomicWFQueueTestFixture, dequeue_enqueue_threads_eq_qsize_2_32) {
+  int num_blocks {2};
+  int block_size {32};
+  int wf_size {this->wf_size};
+  int queue_size {num_blocks * ((block_size - 1) / wf_size + 1)};
+  dequeue_enqueue(num_blocks, block_size, queue_size);
+}
+
+TEST_F(AtomicWFQueueTestFixture, dequeue_enqueue_threads_eq_qsize_32_192) {
+  int num_blocks {32};
+  int block_size {192};
+  int wf_size {this->wf_size};
+  int queue_size {num_blocks * ((block_size - 1) / wf_size + 1)};
+  dequeue_enqueue(num_blocks, block_size, queue_size);
+}
+
+TEST_F(AtomicWFQueueTestFixture, dequeue_enqueue_threads_eq_qsize_64_256) {
+  int num_blocks {64};
+  int block_size {256};
+  int wf_size {this->wf_size};
+  int queue_size {num_blocks * ((block_size - 1) / wf_size + 1)};
+  dequeue_enqueue(num_blocks, block_size, queue_size);
+}
+
+TEST_F(AtomicWFQueueTestFixture, dequeue_enqueue_threads_ge_qsize_16_32_4) {
+  int num_blocks {16};
+  int block_size {32};
+  int queue_size {4};
+  dequeue_enqueue(num_blocks, block_size, queue_size);
+}
+
+TEST_F(AtomicWFQueueTestFixture, dequeue_enqueue_threads_ge_qsize_16_96_8) {
+  int num_blocks {16};
+  int block_size {96};
+  int queue_size {8};
+  dequeue_enqueue(num_blocks, block_size, queue_size);
+}
+
+TEST_F(AtomicWFQueueTestFixture, dequeue_enqueue_threads_ge_qsize_16_320_8) {
+  int num_blocks {16};
+  int block_size {320};
+  int queue_size {8};
+  dequeue_enqueue(num_blocks, block_size, queue_size);
+}
+
+TEST_F(AtomicWFQueueTestFixture, dequeue_enqueue_threads_ge_qsize_32_192_16) {
+  int num_blocks {32};
+  int block_size {192};
+  int queue_size {16};
+  dequeue_enqueue(num_blocks, block_size, queue_size);
+}
+
+TEST_F(AtomicWFQueueTestFixture, dequeue_enqueue_threads_ge_qsize_64_256_64) {
+  int num_blocks {64};
+  int block_size {256};
+  int queue_size {64};
+  dequeue_enqueue(num_blocks, block_size, queue_size);
+}
+
+TEST_F(AtomicWFQueueTestFixture, dequeue_enqueue_threads_ge_qsize_64_576_64) {
+  int num_blocks {64};
+  int block_size {576};
+  int queue_size {64};
+  dequeue_enqueue(num_blocks, block_size, queue_size);
+}
+
+TEST_F(AtomicWFQueueTestFixture, dequeue_enqueue_qsize_ge_threads_4_8_16) {
+  int num_blocks {4};
+  int block_size {8};
+  int queue_size {16};
+  dequeue_enqueue(num_blocks, block_size, queue_size);
+}
+
+TEST_F(AtomicWFQueueTestFixture, dequeue_enqueue_qsize_ge_threads_4_96_16) {
+  int num_blocks {4};
+  int block_size {96};
+  int queue_size {16};
+  dequeue_enqueue(num_blocks, block_size, queue_size);
+}
+
+TEST_F(AtomicWFQueueTestFixture, dequeue_enqueue_qsize_ge_threads_32_192_128) {
+  int num_blocks {32};
+  int block_size {192};
+  int queue_size {128};
+  dequeue_enqueue(num_blocks, block_size, queue_size);
+}
+
+TEST_F(AtomicWFQueueTestFixture, dequeue_enqueue_qsize_ge_threads_32_320_256) {
+  int num_blocks {32};
+  int block_size {320};
+  int queue_size {256};
+  dequeue_enqueue(num_blocks, block_size, queue_size);
+}
+
+TEST_F(AtomicWFQueueTestFixture, dequeue_enqueue_qsize_ge_threads_64_16_128) {
+  int num_blocks {64};
+  int block_size {16};
+  int queue_size {128};
+  dequeue_enqueue(num_blocks, block_size, queue_size);
+}
+
+TEST_F(AtomicWFQueueTestFixture, dequeue_enqueue_qsize_ge_threads_64_96_256) {
+  int num_blocks {64};
+  int block_size {96};
+  int queue_size {256};
+  dequeue_enqueue(num_blocks, block_size, queue_size);
+}
+
+TEST_F(AtomicWFQueueTestFixture, dequeue_enqueue_qsize_ge_threads_64_320_512) {
+  int num_blocks {64};
+  int block_size {320};
+  int queue_size {512};
+  dequeue_enqueue(num_blocks, block_size, queue_size);
+}
+
+TEST_F(AtomicWFQueueTestFixture, dequeue_enqueue_qsize_ge_threads_64_576_1024) {
+  int num_blocks {64};
+  int block_size {576};
+  int queue_size {1024};
+  dequeue_enqueue(num_blocks, block_size, queue_size);
+}
@@ -0,0 +1,170 @@
+/******************************************************************************
+ * 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_ATOMIC_WF_QUEUE_GTEST_HPP
+#define ROCSHMEM_ATOMIC_WF_QUEUE_GTEST_HPP
+
+#include "../src/containers/atomic_wf_queue_impl.hpp"
+#include "gtest/gtest.h"
+#include "../src/memory/hip_allocator.hpp"
+#include "../src/util.hpp"
+#include <iostream>
+
+namespace rocshmem {
+
+template <typename AWFQueue>
+__global__ void wf_lane_ids(AWFQueue* awf_queue, unsigned int* device_array,
+                           int wf_size) {
+
+  int t_id {get_flat_id()};
+  int lane_id {t_id % wf_size};
+  device_array[t_id] = awf_queue->active_logical_lane_id();
+}
+
+template <typename AWFQueue>
+__global__ void concurrent_enqueue_dequeue(
+    AWFQueue* awf_queue,
+    unsigned int* device_array) {
+
+  int t_id {get_flat_id()};
+  int val {awf_queue->dequeue()};
+  device_array[t_id] = val;
+  awf_queue->enqueue(val);
+}
+
+class AtomicWFQueueTestFixture : public ::testing::Test {
+ public:
+  AtomicWFQueueTestFixture() {
+    awf_queue = awf_queue_proxy.get();
+
+    int device_id {};
+    hipDeviceProp_t device_props;
+    CHECK_HIP(hipGetDevice(&device_id));
+    CHECK_HIP(hipGetDeviceProperties(&device_props, device_id));
+    
+    wf_size = device_props.warpSize;
+  }
+
+  ~AtomicWFQueueTestFixture() {}
+
+  void get_thread_lane_ids(unsigned int num_threads = 4) {
+    
+    unsigned int *device_array {nullptr};
+
+    hip_allocator_.allocate(reinterpret_cast<void**>(&device_array),
+                            sizeof(unsigned int) * num_threads);
+  
+
+    hipLaunchKernelGGL(wf_lane_ids, 1, num_threads, 0, nullptr,
+                       awf_queue, device_array, wf_size);
+
+    CHECK_HIP(hipDeviceSynchronize());
+
+    hip_allocator_.deallocate(device_array);
+
+    for (unsigned int i{0}; i < num_threads; i++) {
+      EXPECT_EQ(device_array[i], i % wf_size);
+    }
+  }
+
+  void init_queue(int num_blocks, int block_size) {
+    int num_elems = num_blocks * ((block_size - 1) / wf_size + 1);
+
+    awf_queue->allocate_queue(num_elems);
+
+    for (int i{0}; i < num_elems; i++) {
+      awf_queue->push(i);
+    }
+
+    EXPECT_EQ(awf_queue->get_curr_size(), num_elems);
+    EXPECT_EQ(awf_queue->get_queue_size(), num_elems);
+    EXPECT_EQ(awf_queue->get_tail(), awf_queue->get_head());
+
+    awf_queue->deallocate_queue();
+  }
+
+  void verify(unsigned int *arr, int num_blocks, int block_size) {
+    unsigned int expected_val {};
+    unsigned int lane_id {};
+    unsigned int idx {};
+    for (unsigned int i{0}; i < num_blocks; i++) {
+      for (unsigned int j{0}; j < block_size; j++) {
+        idx = i * block_size + j;
+        lane_id = j % wf_size;
+        if (!lane_id) {
+          expected_val = arr[idx];
+        }
+        EXPECT_EQ(arr[idx], expected_val + lane_id);
+      }
+    }
+  }
+
+  void dequeue_enqueue(int num_blocks, int block_size, int queue_size) {
+
+    int num_threads {num_blocks * block_size};
+    unsigned int *device_array {nullptr};
+
+    hip_allocator_.allocate(reinterpret_cast<void**>(&device_array),
+                            sizeof(unsigned int) * num_threads);
+
+    awf_queue->allocate_queue(queue_size);
+
+    for (int i{0}; i < queue_size; i++) {
+      awf_queue->push(i);
+    }
+
+    EXPECT_EQ(awf_queue->get_curr_size(), queue_size);
+    EXPECT_EQ(awf_queue->get_queue_size(), queue_size);
+    EXPECT_EQ(awf_queue->get_tail(), awf_queue->get_head());
+
+    hipLaunchKernelGGL(concurrent_enqueue_dequeue, num_blocks, block_size, 0,
+                       nullptr, awf_queue, device_array);
+    CHECK_HIP(hipDeviceSynchronize());
+
+    EXPECT_EQ(awf_queue->get_curr_size(), queue_size);
+    EXPECT_EQ(awf_queue->get_tail(), awf_queue->get_head());
+
+    verify(device_array, num_blocks, block_size);
+
+    hip_allocator_.deallocate(device_array);
+    awf_queue->deallocate_queue();
+  }
+
+  protected:
+
+    AtomicWFQueueProxy<HIPAllocator, int> awf_queue_proxy{};
+    AtomicWFQueue<int, HIPAllocator>* awf_queue{};
+
+    /**
+     * @brief An allocator to create objects in device memory.
+     */
+    HIPAllocator hip_allocator_ {};
+
+    /**
+     * @brief Wavefront size.
+     */
+    int wf_size {};
+};
+
+}  // namespace rocshmem
+
+#endif  // ROCSHMEM_ATOMIC_WF_QUEUE_GTEST_HPP