From 0fd628458c2a6eed294e35825147204829598200 Mon Sep 17 00:00:00 2001 From: Brandon Potter Date: Thu, 10 Apr 2025 14:47:24 -0500 Subject: [PATCH] Cleanup unused code in repository (#75) * Remove unused forward_list * Remove unused __read_clock function * Replace wallClk code with hip function * Remove unused unit test for ipc * Remove slab heap * Remove unused EBO spinlock --- src/assembly.hpp | 21 - src/containers/forward_list.hpp | 510 ------------------ src/containers/forward_list_impl.hpp | 233 -------- src/context.hpp | 1 - src/memory/CMakeLists.txt | 1 - src/memory/slab_heap.cpp | 116 ---- src/memory/slab_heap.hpp | 195 ------- src/reverse_offload/backend_ro.cpp | 7 +- src/sync/CMakeLists.txt | 1 - src/sync/spin_ebo_block_mutex.cpp | 131 ----- src/sync/spin_ebo_block_mutex.hpp | 85 --- src/util.cpp | 27 - src/util.hpp | 3 - tests/unit_tests/CMakeLists.txt | 4 - tests/unit_tests/context_ipc_gtest.cpp | 31 -- tests/unit_tests/context_ipc_gtest.hpp | 46 -- tests/unit_tests/forward_list_gtest.cpp | 51 -- tests/unit_tests/forward_list_gtest.hpp | 50 -- tests/unit_tests/slab_heap_gtest.cpp | 130 ----- tests/unit_tests/slab_heap_gtest.hpp | 113 ---- .../unit_tests/spin_ebo_block_mutex_gtest.cpp | 69 --- .../unit_tests/spin_ebo_block_mutex_gtest.hpp | 140 ----- 22 files changed, 6 insertions(+), 1959 deletions(-) delete mode 100644 src/containers/forward_list.hpp delete mode 100644 src/containers/forward_list_impl.hpp delete mode 100644 src/memory/slab_heap.cpp delete mode 100644 src/memory/slab_heap.hpp delete mode 100644 src/sync/spin_ebo_block_mutex.cpp delete mode 100644 src/sync/spin_ebo_block_mutex.hpp delete mode 100644 tests/unit_tests/context_ipc_gtest.cpp delete mode 100644 tests/unit_tests/context_ipc_gtest.hpp delete mode 100644 tests/unit_tests/forward_list_gtest.cpp delete mode 100644 tests/unit_tests/forward_list_gtest.hpp delete mode 100644 tests/unit_tests/slab_heap_gtest.cpp delete mode 100644 tests/unit_tests/slab_heap_gtest.hpp delete mode 100644 tests/unit_tests/spin_ebo_block_mutex_gtest.cpp delete mode 100644 tests/unit_tests/spin_ebo_block_mutex_gtest.hpp diff --git a/src/assembly.hpp b/src/assembly.hpp index 72720ef75f..fa27e24c36 100644 --- a/src/assembly.hpp +++ b/src/assembly.hpp @@ -245,27 +245,6 @@ __device__ __forceinline__ void store_asm(uint8_t* val, uint8_t* dst, } } -__device__ __forceinline__ uint64_t __read_clock() { - uint64_t clock{}; -#if defined(__gfx906__) -#endif -#if defined(__gfx908__) -#endif -#if defined(__gfx90a__) - asm volatile( - "s_memrealtime %0\n" - "s_waitcnt lgkmcnt(0)\n" - : "=s"(clock)); -#endif -#if defined(__gfx942__) - asm volatile( - "s_memrealtime %0\n" - "s_waitcnt lgkmcnt(0)\n" - : "=s"(clock)); -#endif - return clock; -} - } // namespace rocshmem #endif // LIBRARY_SRC_ASSEMBLY_HPP_ diff --git a/src/containers/forward_list.hpp b/src/containers/forward_list.hpp deleted file mode 100644 index acd0d70343..0000000000 --- a/src/containers/forward_list.hpp +++ /dev/null @@ -1,510 +0,0 @@ -/****************************************************************************** - * 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 LIBRARY_SRC_CONTAINERS_FORWARD_LIST_HPP_ -#define LIBRARY_SRC_CONTAINERS_FORWARD_LIST_HPP_ - -#include - -#include -#include -#include - -#include "../src/device_proxy.hpp" -#include "../src/memory/hip_allocator.hpp" -#include "../sync/abql_block_mutex.hpp" - -namespace rocshmem { - -/***************************************************************************** - ******************************* FORWARD LIST ******************************** - *****************************************************************************/ - -template -class ForwardList { - friend class ForwardListTestFixture; - - struct Node { - TYPE data; - Node* next; - }; - - template - class Iterator; - - public: - using iterator = Iterator; - using const_iterator = Iterator; - - /** - * @brief Default constructor - */ - ForwardList() = default; - - explicit ForwardList(const ALLOC& alloc); - - /** - * @brief Fill constructor - */ - explicit ForwardList(size_t n, const ALLOC& alloc = ALLOC()); - - explicit ForwardList(size_t n, const TYPE& val, const ALLOC& alloc = ALLOC()); - - /** - * @brief Range constructor - */ - template - ForwardList(InputIterator first, InputIterator last, - const ALLOC& alloc = ALLOC()); - - /** - * @brief Copy constructor - */ - ForwardList(const ForwardList& fwdlst); - - ForwardList(const ForwardList& fwdlst, const ALLOC& alloc); - - /** - * @brief Move constructor - */ - ForwardList(ForwardList&& fwdlst); - - ForwardList(ForwardList&& fwdlst, const ALLOC& alloc); - - /** - * @brief Initializer list constructor - */ - ForwardList(std::initializer_list il, const ALLOC& alloc = ALLOC()); - - /** - * @brief Destructor - */ - ~ForwardList(); - - /** - * @brief Copy assignment operator - */ - ForwardList& operator=(const ForwardList& fwdlst); - - /** - * @brief Move assignment operator - */ - ForwardList& operator=(ForwardList&& fwdlst); - - /** - * @brief Initializer list assignment - */ - ForwardList& operator=(std::initializer_list il); - - /** - * @brief Returns iterator pointing to position before first element. - * - * The iterator returned shall not be dereferenced: - * It is meant to be used as an argument for member functions - * emplace_after, insert_after, erase_after or splice_after, to specify - * the beginning of the sequence as the location where the action is - * performed. - */ - iterator before_begin() noexcept; - - /** - * @brief Returns iterator pointing to first element in ForwardList. - * - * Notice that, unlike member ForwardList::front, which returns a - * reference to the first element, this function returns a forward iterator - * pointing to it. - * - * If the container is empty, the returned iterator value shall not be - * dereferenced. - */ - iterator begin() noexcept; - - const_iterator begin() const noexcept; - - /** - * @brief Returns an iterator to past-the-end element in ForwardList. - */ - iterator end() noexcept; - - const_iterator end() const noexcept; - - /** - * @brief Returns const_iterator pointing to position before first element. - */ - const_iterator cbefore_begin() const noexcept; - - /** - * @brief Returns const_iterator pointing to first element. - */ - const_iterator cbegin() const noexcept; - - /** - * @brief Returns const_iterator pointing to past-the-end element. - */ - const_iterator cend() const noexcept; - - /** - * @brief Returns bool indicating if ForwardList container is empty. - */ - bool empty() const noexcept; - - /** - * @brief Returns maximum number of elements that ForwardList can hold. - */ - size_t max_size() const noexcept; - - /** - * @brief Returns reference to first element in the ForwardList. - */ - ForwardList& front(); - - const ForwardList& front() const; - - /** - * @brief Assigns new values, replacing current contents, and modify size. - * - * In range version, the new contents are elements constructed from each - * of the elements in the range between first and last, in the same order. - */ - template - void assign(InputIterator first, InputIterator last); - - /** - * @brief Assigns new values, replacing current contents, and modify size. - * - * In fill version, the new contents are n elements, each initialized to - * a copy of val. - */ - void assign(size_t n, const TYPE& val); - - /** - * @brief Assigns new values, replacing current contents, and modify size. - * - * In initializer list version, the new contents are copies of the values - * passed as initializer list, in the same order. - */ - void assign(std::initializer_list il); - - /** - * @brief Inserts new element at beginning of the ForwardList. - * - * The element goes into the container right before its current first - * element. This new element is constructed in place using args as the - * arguments for its construction. - */ - template - void emplace_front(Args&&... args); - - /** - * @brief Inserts new element at the beginning of the ForwardList. - * - * The element goes into the container right before its current first - * element. The content of val is copied (or moved) to the inserted - * element. - */ - void push_front(const TYPE& val); - - void push_front(TYPE&& val); - - /** - * @brief Removes first element in ForwardList, reducing its size by one. - */ - void pop_front(); - - /** - * @brief Inserts a new element after element at position. - * - * This new element is constructed in place using args as the arguments - * for its construction. - */ - template - iterator emplace_after(const_iterator position, Args&&... args); - - /** - * @brief Inserts new elements after the element at position. - */ - iterator insert_after(const_iterator position, const TYPE& val); - - iterator insert_after(const_iterator position, TYPE&& val); - - iterator insert_after(const_iterator position, size_t n, const TYPE& val); - - template - iterator insert_after(const_iterator position, InputIterator first, - InputIterator last); - - iterator insert_after(const_iterator position, - std::initializer_list il); - - /** - * @brief Removes either a single element or a range of elements. - */ - iterator erase_after(const_iterator position); - - iterator erase_after(const_iterator position, const_iterator last); - - /** - * @brief Exchanges content by the content of fwdlst. - * - * Sizes may differ. - */ - void swap(ForwardList& fwdlst); - - /** - * @brief Resizes the container to contain n elements. - */ - void resize(size_t n); - - void resize(size_t n, const TYPE& val); - - /** - * @brief Removes all elements, all leaves container with size 0. - */ - void clear() noexcept; - - /** - * @brief Transfers all elements of fwdlist into container. - */ - void splice_after(const_iterator position, ForwardList& fwdlst); // NOLINT - - void splice_after(const_iterator position, ForwardList&& fwdlst); - - /** - * @brief Transfers only elements pointed by from fwdlist into container. - */ - void splice_after(const_iterator position, ForwardList& fwdlst, // NOLINT - const_iterator i); - - void splice_after(const_iterator position, ForwardList&& fwdlst, - const_iterator i); - - /** - * @brief Transfers the range (first,last) from fwdlist into container. - */ - void splice_after(const_iterator position, ForwardList& fwdlst, // NOLINT - const_iterator first, const_iterator last); - - void splice_after(const_iterator position, ForwardList&& fwdlst, // NOLINT - const_iterator first, const_iterator last); - - /** - * @brief Removes all the elements that compare equal to val. - */ - void remove(const TYPE& val); - - /** - * @brief Removes all elements for which Predicate pred returns true. - */ - template - void remove_if(Predicate pred); - - /** - * @brief Removes all but first element from consecutive group equal elems. - */ - void unique(); - - /** - * @brief Takes comparison func that determines "uniqueness" of elem. - */ - template - void unique(BinaryPredicate binary_pred); - - /** - * @brief Merges ForwardList by transferring all elements. - * - * Both containers shall already be ordered before calling merge. - */ - void merge(ForwardList& fwdlst); // NOLINT - - void merge(ForwardList&& fwdlst); // NOLINT - - /** - * @brief Same as merge, but take specific predicate to perform comparison. - */ - template - void merge(ForwardList& fwdlst, Compare comp); // NOLINT - - template - void merge(ForwardList&& fwdlst, Compare comp); // NOLINT - - /** - * @brief Sorts elements in ForwardList, altering position in container. - */ - void sort(); - - template - void sort(Compare comp); // NOLINT - - /** - * @brief Reverses order of elements in the ForwardList container. - */ - void reverse() noexcept; - - /** - * @brief Returns a copy of the allocator object associated with container. - */ - ALLOC - get_allocator() const noexcept; - - private: - /** - * @brief Internal memory allocator used to create list nodes. - */ - MemoryAllocator allocator_{}; - - /** - * @brief First element in the list. - */ - Node* head_{nullptr}; - - /** - * @brief Last element in the list. - */ - Node* tail_{nullptr}; - - /** - * @brief Size of the list. - */ - size_t size_{0}; -}; - -/***************************************************************************** - ********************************* ITERATOR ********************************** - *****************************************************************************/ - -template -template -class ForwardList::Iterator { - using NodeT = typename ForwardList::Node; - - public: - using iterator_category = std::forward_iterator_tag; - using difference_type = std::ptrdiff_t; - using value_type = TYPE; - using pointer = TYPE*; - using reference = TYPE&; - - Iterator(NodeT* ptr); - - Iterator operator++(); - - Iterator operator++(int); - - template - typename std::enable_if::type operator*(); - - template - typename std::enable_if::type operator*(); - - pointer operator->(); - - template - friend bool operator==(ITER_TYPE& a, ITER_TYPE& b); // NOLINT - - template - friend bool operator==(ITER_TYPE& a, ITER_TYPE b); // NOLINT - - template - friend bool operator!=(ITER_TYPE& a, ITER_TYPE& b); // NOLINT - - template - friend bool operator!=(ITER_TYPE& a, ITER_TYPE b); // NOLINT - - private: - NodeT* node_ptr_; -}; - -/***************************************************************************** - ******************************* COMPARATORS ********************************* - *****************************************************************************/ - -template -bool operator==(const ForwardList& lhs, - const ForwardList& rhs); - -template -bool operator==(const ForwardList& lhs, const std::string rhs); - -template -bool operator==(const std::string lhs, const ForwardList& rhs); - -template -bool operator!=(const ForwardList& lhs, - const ForwardList& rhs); - -template -bool operator<(const ForwardList& lhs, - const ForwardList& rhs); - -template -bool operator<=(const ForwardList& lhs, - const ForwardList& rhs); - -template -bool operator>(const ForwardList& lhs, - const ForwardList& rhs); - -template -bool operator>=(const ForwardList& lhs, - const ForwardList& rhs); - -template -void swap(ForwardList& x, ForwardList& y); // NOLINT - -/***************************************************************************** - ******************************** STRINGIFIERS ******************************* - *****************************************************************************/ - -template -std::ostream& operator<<(std::ostream& os, ForwardList const& list); - -template -std::string to_string(const ForwardList& list); - -/***************************************************************************** - *********************************** PROXY *********************************** - *****************************************************************************/ - -template -class ForwardListProxy { - using ProxyT = DeviceProxy>; - - public: - ForwardList(size_t num_elems = 1) : proxy_{num_elems} {} - - ForwardList(const ForwardList& other) = delete; - - ForwardList& operator=(const ForwardList& other) = delete; - - ForwardList(ForwardList&& other) = default; - - ForwardList& operator=(ForwardList&& other) = default; - - __host__ __device__ ForwardList* get() { return proxy_.get(); } - - private: - ProxyT proxy_{}; -}; - -} // namespace rocshmem - -#endif // LIBRARY_SRC_CONTAINERS_FORWARD_LIST_HPP_ diff --git a/src/containers/forward_list_impl.hpp b/src/containers/forward_list_impl.hpp deleted file mode 100644 index 7e354aa9b9..0000000000 --- a/src/containers/forward_list_impl.hpp +++ /dev/null @@ -1,233 +0,0 @@ -/****************************************************************************** - * 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 LIBRARY_SRC_CONTAINERS_FORWARD_LIST_IMPL_HPP_ -#define LIBRARY_SRC_CONTAINERS_FORWARD_LIST_IMPL_HPP_ - -#include -#include -#include - -#include "forward_list.hpp" - -namespace rocshmem { - -/***************************************************************************** - ******************************* FORWARD LIST ******************************** - *****************************************************************************/ - -template -ForwardList::~ForwardList() { - while (head_ != nullptr) { - Node* temp{head_}; - head_ = head_->next; - allocator_.deallocate(temp); - } - tail_ = nullptr; -} - -template -ForwardList::ForwardList(const ForwardList& fwdlst) {} - -template -ForwardList::ForwardList(const ForwardList& fwdlst, - const ALLOC& alloc) {} - -template -ForwardList::ForwardList(std::initializer_list il, - const ALLOC& allocator) { - allocator_ = allocator; - assign(il); -} - -template -typename ForwardList::iterator -ForwardList::begin() noexcept { - iterator i(head_); - return i; -} - -template -typename ForwardList::const_iterator -ForwardList::begin() const noexcept { - const_iterator i(head_); - return i; -} - -template -typename ForwardList::iterator -ForwardList::end() noexcept { - iterator i(tail_); - return i; -} - -template -typename ForwardList::const_iterator -ForwardList::end() const noexcept { - const_iterator i(tail_); - return i; -} - -template -void ForwardList::assign(std::initializer_list il) { - resize(il.size()); - std::copy_n(il.begin(), il.size(), begin()); -} - -template -void ForwardList::resize(size_t n) {} - -template -void ForwardList::resize(size_t n, const TYPE& val) {} - -template -void ForwardList::clear() noexcept {} - -/***************************************************************************** - ********************************* ITERATOR ********************************** - *****************************************************************************/ - -template -template -ForwardList::Iterator::Iterator(NodeT* ptr) - : node_ptr_(ptr) {} - -template -template -typename ForwardList::template Iterator -ForwardList::Iterator::operator++() { - if (node_ptr_) { - node_ptr_ = node_ptr_->next; - } - return *this; -} - -template -template -typename ForwardList::template Iterator -ForwardList::Iterator::operator++(int) { - Iterator iterator = *this; - ++*this; - return iterator; -} - -template -template -template -typename std::enable_if::type -ForwardList::Iterator::operator*() { - return node_ptr_->data; -} - -template -template -template -typename std::enable_if::type -ForwardList::Iterator::operator*() { - return node_ptr_->data; -} - -template -template -TYPE* ForwardList::Iterator::operator->() { - return &node_ptr_->data; -} - -template -bool operator==(ITER_TYPE& a, ITER_TYPE& b) { // NOLINT - return a.node_ptr_ == b.node_ptr_; -} - -template -bool operator==(ITER_TYPE& a, ITER_TYPE b) { // NOLINT - return a.node_ptr_ == b.node_ptr_; -} - -template -bool operator!=(ITER_TYPE& a, ITER_TYPE& b) { // NOLINT - return a.node_ptr_ != b.node_ptr_; -} - -template -bool operator!=(ITER_TYPE& a, ITER_TYPE b) { // NOLINT - return a.node_ptr_ != b.node_ptr_; -} - -/***************************************************************************** - ******************************* COMPARATORS ********************************* - *****************************************************************************/ - -template -bool operator==(const ForwardList& lhs, - const ForwardList& rhs) { - return false; -} - -template -bool operator==(const ForwardList& lhs, const std::string rhs) { - return false; -} - -template -bool operator==(const std::string lhs, const ForwardList& rhs) { - return false; -} - -template -bool operator!=(const ForwardList& lhs, - const ForwardList& rhs) { - return false; -} - -template -bool operator!=(const ForwardList& lhs, const std::string rhs) { - return false; -} - -template -bool operator!=(const std::string lhs, const ForwardList& rhs) { - return false; -} - -/***************************************************************************** - ******************************** STRINGIFIERS ******************************* - *****************************************************************************/ - -template -std::ostream& operator<<(std::ostream& os, ForwardList const& list) { - using CItr = typename ForwardList::const_iterator; - for (CItr ci = list.begin(); ci != list.end(); ++ci) { - os << *ci << " "; - } - return os; -} - -template -std::string to_string(const ForwardList& list) { - std::stringstream ss; - ss << list; - return ss.str(); -} - -} // namespace rocshmem - -#endif // LIBRARY_SRC_CONTAINERS_FORWARD_LIST_IMPL_HPP_ diff --git a/src/context.hpp b/src/context.hpp index e74299fe0c..2badbd2ba9 100644 --- a/src/context.hpp +++ b/src/context.hpp @@ -30,7 +30,6 @@ #include "host/host.hpp" #include "ipc_policy.hpp" #include "stats.hpp" -#include "sync/spin_ebo_block_mutex.hpp" #include "wf_coal_policy.hpp" namespace rocshmem { diff --git a/src/memory/CMakeLists.txt b/src/memory/CMakeLists.txt index cf72f0fc2f..a481e99372 100644 --- a/src/memory/CMakeLists.txt +++ b/src/memory/CMakeLists.txt @@ -27,6 +27,5 @@ target_sources( ${PROJECT_NAME} PRIVATE single_heap.cpp - slab_heap.cpp memory_allocator.cpp ) diff --git a/src/memory/slab_heap.cpp b/src/memory/slab_heap.cpp deleted file mode 100644 index c126111770..0000000000 --- a/src/memory/slab_heap.cpp +++ /dev/null @@ -1,116 +0,0 @@ -/****************************************************************************** - * 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 "slab_heap.hpp" - -#include - -#include "../util.hpp" - -namespace rocshmem { - -SlabHeap::SlabHeap() { - if (auto slab_size_cstr = getenv("ROCSHMEM_SLAB_SIZE")) { - std::stringstream sstream(slab_size_cstr); - size_t slab_size; - sstream >> slab_size; - heap_mem_ = HEAP_T{slab_size}; - strat_ = STRAT_T{&heap_mem_}; - } -} - -void SlabHeap::malloc(void** ptr, size_t size) { - strat_.alloc(reinterpret_cast(ptr), size); -} - -__device__ void SlabHeap::malloc(void** ptr, size_t size) { - /* - * Grab the mutex from the proxy object which owns it. - */ - auto mutex{mutex_.get()}; - - /* - * Take the ticketed lock. - * - * The lock is held jointly by all threads in the block. - */ - auto ticket{mutex->lock()}; - - /* - * Perform allocation and verify it worked. - * - * Allocation should only be run by only one thread in the - * strategy code. - */ - char** ptr_c{reinterpret_cast(ptr)}; - strat_.alloc(ptr_c, size); - __threadfence(); - - /* - * The notifier works with uint64_t for the address broadcasts - * between threads (as type erasure for the pointer arithmetic). - */ - uint64_t ptr_deref_u64{reinterpret_cast(*ptr)}; - - /* - * Notify other threads in block about the allocation result. - */ - auto notifier{notifier_.get()}; - if (!threadIdx.x) { - notifier->store(ptr_deref_u64); - notifier->fence(); - } - __syncthreads(); - uint64_t notification_u64{notifier->load()}; - - /* - * Write to the ptr parameter (to return it back up the call stack). - */ - char* read_value_c{reinterpret_cast(notification_u64)}; - *ptr_c = read_value_c; - - /* - * Release the lock with our ticket number. - */ - mutex->unlock(ticket); -} - -__host__ __device__ void SlabHeap::free(void* ptr) { - if (!ptr) { - return; - } - strat_.free(reinterpret_cast(ptr)); -} - -void* SlabHeap::realloc(void* ptr, size_t size) { return nullptr; } - -void* SlabHeap::malign(size_t alignment, size_t size) { return nullptr; } - -char* SlabHeap::get_base_ptr() { return heap_mem_.get_ptr(); } - -size_t SlabHeap::get_size() { return heap_mem_.get_size(); } - -size_t SlabHeap::get_used() { return strat_.current() - get_base_ptr(); } - -size_t SlabHeap::get_avail() { return get_size() - get_used(); } - -} // namespace rocshmem diff --git a/src/memory/slab_heap.hpp b/src/memory/slab_heap.hpp deleted file mode 100644 index 23bd4e099f..0000000000 --- a/src/memory/slab_heap.hpp +++ /dev/null @@ -1,195 +0,0 @@ -/****************************************************************************** - * 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 LIBRARY_SRC_MEMORY_SLAB_HEAP_HPP_ -#define LIBRARY_SRC_MEMORY_SLAB_HEAP_HPP_ - -#include "dev_mono_linear.hpp" -#include "heap_memory.hpp" -#include "heap_type.hpp" -#include "notifier.hpp" -#include "../sync/abql_block_mutex.hpp" - -/** - * @file slab_heap.hpp - * - * @brief Contains a heap used to allocate library objects - * - * The slab heap is used internally by the library - */ - -namespace rocshmem { - -class SlabHeap { - /** - * @brief Helper type for allocation strategy - */ - using STRAT_T = DevMonoLinear; - - /** - * @brief Helper type for notifier - */ - using NOTIFIER_PROXY_T = NotifierProxy; - - /** - * @brief Helper type for notifier - */ - using MUTEX_PROXY_T = ABQLBlockMutexProxy; - - public: - /** - * @brief Primary constructor - */ - SlabHeap(); - - /** - * @brief Allocates memory from the heap - * - * @param[in,out] A pointer to memory handle - * @param[in] Size in bytes of memory allocation - */ - void malloc(void** ptr, size_t size); - - /** - * @brief Allocates memory from the heap - * - * @param[in,out] A pointer to memory handle - * @param[in] Size in bytes of memory allocation - */ - __device__ void malloc(void** ptr, size_t size); - - /** - * @brief Frees memory from the heap - * - * @param[in] Raw pointer to heap memory - */ - __host__ __device__ void free(void* ptr); - - /** - * @brief - * - * @param[in] - * @param[in] - * - * @return - */ - void* realloc(void* ptr, size_t size); - - /** - * @brief - * - * @param[in] - * @param[in] - * - * @return - */ - void* malign(size_t alignment, size_t size); - - /** - * @brief Accessor for heap base ptr - * - * @return Pointer to base of my heap - */ - char* get_base_ptr(); - - /** - * @brief Accessor for heap size - * - * @return Amount of bytes in heap - */ - size_t get_size(); - - /** - * @brief Accessor for heap usage - * - * @return Amount of used bytes in heap - */ - size_t get_used(); - - /** - * @brief Accessor for heap available - * - * @return Amount of available bytes in heap - */ - size_t get_avail(); - - private: - /** - * @brief Heap memory object - */ - HEAP_T heap_mem_{}; - - /** - * @brief Allocation strategy object - */ - STRAT_T strat_{&heap_mem_}; - - /** - * @brief Notifier proxy to share information between threads. - * - * Need this object to share allocation information between the - * leader thread (that does allocation) and the follower threads - * (who need the allocation address). - */ - NOTIFIER_PROXY_T notifier_{}; - - /** - * @brief Mutex to access the heap mutator methods. - */ - MUTEX_PROXY_T mutex_; -}; - -template -class SlabHeapProxy { - using ProxyT = DeviceProxy; - - public: - /* - * Placement new the memory which is allocated by proxy_ - */ - SlabHeapProxy(size_t num_elems = 1) : proxy_{num_elems} { - new (proxy_.get()) SlabHeap(); - } - - SlabHeapProxy(const SlabHeapProxy& other) = delete; - - SlabHeapProxy& operator=(const SlabHeapProxy& other) = delete; - - SlabHeapProxy(SlabHeapProxy&& other) = default; - - SlabHeapProxy& operator=(SlabHeapProxy&& other) = default; - - /* - * Since placement new is called in the constructor, then - * delete must be called manually. - */ - ~SlabHeapProxy() { proxy_.get()->~SlabHeap(); } - - __host__ __device__ SlabHeap* get() { return proxy_.get(); } - - private: - ProxyT proxy_{}; -}; - -} // namespace rocshmem - -#endif // LIBRARY_SRC_MEMORY_SLAB_HEAP_HPP_ diff --git a/src/reverse_offload/backend_ro.cpp b/src/reverse_offload/backend_ro.cpp index bdcd54df98..f2d4b9c068 100644 --- a/src/reverse_offload/backend_ro.cpp +++ b/src/reverse_offload/backend_ro.cpp @@ -242,7 +242,12 @@ void ROBackend::dump_backend_stats() { total += globalStats.getStat(i); } - uint64_t gpu_frequency_mhz{wallClk_freq_mhz()}; + int device_id; + hipDeviceProp_t device_props; + CHECK_HIP(hipGetDevice(&device_id)); + int wallClockMhz; + CHECK_HIP(hipDeviceGetAttribute(&wallClockMhz, hipDeviceAttributeWallClockRate, device_id)); + int gpu_frequency_mhz{wallClockMhz}; uint64_t us_wait_slot{0}; uint64_t us_pack{0}; diff --git a/src/sync/CMakeLists.txt b/src/sync/CMakeLists.txt index 88f2d70840..22e48cf510 100644 --- a/src/sync/CMakeLists.txt +++ b/src/sync/CMakeLists.txt @@ -26,6 +26,5 @@ target_sources( ${PROJECT_NAME} PRIVATE - spin_ebo_block_mutex.cpp abql_block_mutex.cpp ) diff --git a/src/sync/spin_ebo_block_mutex.cpp b/src/sync/spin_ebo_block_mutex.cpp deleted file mode 100644 index 99ba441f80..0000000000 --- a/src/sync/spin_ebo_block_mutex.cpp +++ /dev/null @@ -1,131 +0,0 @@ -/****************************************************************************** - * 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 "../sync/spin_ebo_block_mutex.hpp" - -#include - -#include "../util.hpp" - -namespace rocshmem { - -__device__ __host__ SpinEBOBlockMutex::SpinEBOBlockMutex(bool shareable) - : shareable_(shareable) {} - -__device__ void SpinEBOBlockMutex::lock() { -#ifdef USE_SHARED_CTX - if (!shareable_) { - return; - } - /* - * We need to check this context out to a work-group, and only let threads - * that are a part of the owning work-group through. It's a bit like a - * re-entrant lock, with the added twist that a thread checks out the lock - * for his entire work-group. - */ - int num_threads_in_wv = wave_SZ(); - - if (get_flat_block_id() % WF_SIZE == lowerID()) { - /* - * All the metadata associated with this lock needs to be accessed - * atomically or it will race. - */ - while (atomicCAS(reinterpret_cast(&ctx_lock_), 0, 1) == 1) { - uint64_t time_now = clock64(); - int64_t wait_time = 100; - while (time_now + wait_time > clock64()) { - __threadfence(); - } - wait_time *= 2; - wait_time = min(wait_time, 20000); - } - - /* - * If somebody in my work-group already owns the default context, just - * record how many threads are going to be here and go about our - * business. - * - * If my work-group doesn't own the default context, then - * we need to wait for it to become available. Relinquish - * ctx_lock while waiting or it will never become available. - * - */ - int wg_id = get_flat_grid_id(); - while (wg_owner_ != wg_id) { - if (wg_owner_ == -1) { - wg_owner_ = wg_id; - __threadfence(); - } else { - ctx_lock_ = 0; - __threadfence(); - // Performance is terrible. Backoff slightly helps. - while (atomicCAS(reiterpret_cast(&ctx_lock_), 0, 1) == 1) { - uint64_t time_now = clock64(); - int64_t wait_time = 100; - while (time_now + wait_time > clock64()) { - __threadfence(); - } - wait_time *= 2; - wait_time = min(wait_time, 20000); - } - } - } - - num_threads_in_lock_ += num_threads_in_wv; - __threadfence(); - - ctx_lock_ = 0; - __threadfence(); - } -#endif -} - -__device__ void SpinEBOBlockMutex::unlock() { -#ifdef USE_SHARED_CTX - __threadfence(); - if (!shareable_) { - return; - } - int num_threads_in_wv{wave_SZ()}; - - if (get_flat_block_id() % WF_SIZE == lowerID()) { - while (atomicCAS(reinterpret_cast(&ctx_lock_), 0, 1) == 1) { - } - - num_threads_in_lock_ -= num_threads_in_wv; - - /* - * Last thread out for this work-group opens the door for other - * work-groups to take possession. - */ - if (num_threads_in_lock_ == 0) { - wg_owner_ = -1; - } - - __threadfence(); - - ctx_lock_ = 0; - __threadfence(); - } -#endif -} -} // namespace rocshmem diff --git a/src/sync/spin_ebo_block_mutex.hpp b/src/sync/spin_ebo_block_mutex.hpp deleted file mode 100644 index ed52981c50..0000000000 --- a/src/sync/spin_ebo_block_mutex.hpp +++ /dev/null @@ -1,85 +0,0 @@ -/****************************************************************************** - * 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 LIBRARY_SRC_SYNC_SPIN_EBO_BLOCK_MUTEX_HPP_ -#define LIBRARY_SRC_SYNC_SPIN_EBO_BLOCK_MUTEX_HPP_ - -#include - -namespace rocshmem { - -class SpinEBOBlockMutex { - public: - /** - * @brief Secondary constructor - */ - SpinEBOBlockMutex() = default; - - /** - * @brief Primary constructor - */ - __device__ __host__ SpinEBOBlockMutex(bool shareable); - - /** - * @brief locks the device mutex - * - * @return void - */ - __device__ void lock(); - - /** - * @brief unlocks the device mutex - * - * @return void - */ - __device__ void unlock(); - - /** - * @brief Tells caller if mutex is enabled. - */ - __host__ __device__ bool enabled() { return shareable_; } - - /** - * @brief Context can be shared between different workgroups. - */ - bool shareable_{false}; - - private: - /** - * @brief Shareable context lock. - */ - volatile int ctx_lock_{0}; - - /** - * @brief Shareable context owner. - */ - volatile int wg_owner_{-1}; - - /** - * @brief Number of threads in the owning block inside of locked calls. - */ - volatile int num_threads_in_lock_{0}; -}; - -} // namespace rocshmem - -#endif // LIBRARY_SRC_SYNC_SPIN_EBO_BLOCK_MUTEX_HPP_ diff --git a/src/util.cpp b/src/util.cpp index 337f33dc04..71890a0748 100644 --- a/src/util.cpp +++ b/src/util.cpp @@ -119,33 +119,6 @@ void rocm_memory_lock_to_fine_grain(void* ptr, size_t size, void** gpu_ptr, } } -// TODO(kpunniya): use runtime value instead of hard-coded value -uint64_t wallClk_freq_mhz() { - hipDeviceProp_t deviceProp{}; - CHECK_HIP(hipGetDeviceProperties(&deviceProp, 0)); -// In rocm 6.x and gcnArchName is used instead of gcnArch -#if HIP_VERSION_MAJOR >= 6 - const auto gcnArch = std::atoi(deviceProp.gcnArchName); -#else - const auto gcnArch = deviceProp.gcnArch; -#endif - - switch (gcnArch) { - case 900: // MI25 - return 27; - case 906: - return 25; // MI50,MI60 - case 908: - return 25; // MI100 - case 910: - return 25; // MI200 - default: - assert(false && "clock data unavailable"); - return 0; - } - return 0; -} - struct rocshmem_env_config_t rocshmem_env_config; void rocshmem_env_config_init(void) { diff --git a/src/util.hpp b/src/util.hpp index fcf3e9e496..497730d6b3 100644 --- a/src/util.hpp +++ b/src/util.hpp @@ -266,9 +266,6 @@ int rocm_init(); void rocm_memory_lock_to_fine_grain(void* ptr, size_t size, void** gpu_ptr, int gpu_id); -// Returns clock frequency used by s_memrealtime() in Mhz -uint64_t wallClk_freq_mhz(); - struct rocshmem_env_config_t { int ro_disable_ipc = 0; int ro_progress_delay = 3; diff --git a/tests/unit_tests/CMakeLists.txt b/tests/unit_tests/CMakeLists.txt index e32fde6866..6d25509848 100644 --- a/tests/unit_tests/CMakeLists.txt +++ b/tests/unit_tests/CMakeLists.txt @@ -80,17 +80,13 @@ target_sources( address_record_gtest.cpp index_strategy_gtest.cpp single_heap_gtest.cpp - #slab_heap_gtest.cpp # Test is disabled because class unused symmetric_heap_gtest.cpp pow2_bins_gtest.cpp remote_heap_info_gtest.cpp mpi_init_singleton_gtest.cpp - #spin_ebo_block_mutex_gtest.cpp abql_block_mutex_gtest.cpp notifier_gtest.cpp - #forward_list_gtest.cpp free_list_gtest.cpp - #context_ipc_gtest.cpp wavefront_size_gtest.cpp atomic_wf_queue_gtest.cpp ) diff --git a/tests/unit_tests/context_ipc_gtest.cpp b/tests/unit_tests/context_ipc_gtest.cpp deleted file mode 100644 index c963a65673..0000000000 --- a/tests/unit_tests/context_ipc_gtest.cpp +++ /dev/null @@ -1,31 +0,0 @@ -/****************************************************************************** - * 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 "context_ipc_gtest.hpp" - -using namespace rocshmem; - -TEST_F(ContextIpcTestFixture, constructor) { - /* do nothing for the moment, I *think* the - ** constructor is invoked automatically - */ -} diff --git a/tests/unit_tests/context_ipc_gtest.hpp b/tests/unit_tests/context_ipc_gtest.hpp deleted file mode 100644 index 24a80fcedb..0000000000 --- a/tests/unit_tests/context_ipc_gtest.hpp +++ /dev/null @@ -1,46 +0,0 @@ -/****************************************************************************** - * 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_CONTEXT_IPC_GTEST_HPP -#define ROCSHMEM_CONTEXT_IPC_GTEST_HPP - -#include "gtest/gtest.h" - -#include "../src/ipc/context_ipc_device.hpp" -#include "../src/ipc/backend_ipc.hpp" - -namespace rocshmem { - -class ContextIpcTestFixture : public ::testing::Test -{ - protected: - /** - * @brief Context Ipc Test - */ - IPCBackend be{MPI_COMM_WORLD}; - - IPCContext ipc_context_ {&be}; -}; - -} // namespace rocshmem - -#endif // ROCSHMEM_CONTEXT_IPC_GTEST_HPP diff --git a/tests/unit_tests/forward_list_gtest.cpp b/tests/unit_tests/forward_list_gtest.cpp deleted file mode 100644 index 0c96e77644..0000000000 --- a/tests/unit_tests/forward_list_gtest.cpp +++ /dev/null @@ -1,51 +0,0 @@ -/****************************************************************************** - * 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 "forward_list_gtest.hpp" - -#include -#include - -using namespace rocshmem; - -/***************************************************************************** - ******************************* Fixture Tests ******************************* - *****************************************************************************/ - -TEST_F(ForwardListTestFixture, default_constructor) { - default_constructor_test(); -} - -TEST(ForwardListTest, constructor_tests) { - ForwardList list_1{"rocshmem", "forward_list"}; - std::string str_1 = "rocshmem forward_list"; - ASSERT_EQ(to_string(list_1), str_1); - - // ForwardList list_2(list_1.begin(), list_1.end()); - // std::cout << "list_2: " << list_2 << '\n'; - - // ForwardList list_3(list_1); - // std::cout << "list_3: " << list_3 << '\n'; - - // ForwardList list_4(5, "rocm"); - // std::cout << "list_4: " << list_4 << '\n'; -} diff --git a/tests/unit_tests/forward_list_gtest.hpp b/tests/unit_tests/forward_list_gtest.hpp deleted file mode 100644 index 8c3cf7924b..0000000000 --- a/tests/unit_tests/forward_list_gtest.hpp +++ /dev/null @@ -1,50 +0,0 @@ -/****************************************************************************** - * 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_FORWARD_LIST_GTEST_HPP -#define ROCSHMEM_FORWARD_LIST_GTEST_HPP - -#include "gtest/gtest.h" - -#include "../src/containers/forward_list_impl.hpp" -#include "../src/memory/hip_allocator.hpp" - -namespace rocshmem { - -class ForwardListTestFixture : public ::testing::Test { - public: - ForwardListTestFixture() { - } - - void - default_constructor_test() { - ASSERT_EQ(list_.head_, nullptr); - ASSERT_EQ(to_string(list_), std::string{}); - } - - private: - ForwardList list_ {}; -}; - -} // namespace rocshmem - -#endif // ROCSHMEM_FORWARD_LIST_GTEST_HPP diff --git a/tests/unit_tests/slab_heap_gtest.cpp b/tests/unit_tests/slab_heap_gtest.cpp deleted file mode 100644 index 06e6e82129..0000000000 --- a/tests/unit_tests/slab_heap_gtest.cpp +++ /dev/null @@ -1,130 +0,0 @@ -/****************************************************************************** - * 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 "slab_heap_gtest.hpp" - -using namespace rocshmem; - -TEST_F(SlabHeapTestFixture, malloc_free) { - void *ptr{nullptr}; - size_t request_bytes{48}; - - auto slab{slab_.get()}; - slab->malloc(&ptr, request_bytes); - - ASSERT_NE(ptr, nullptr); - ASSERT_NO_FATAL_FAILURE(slab->free(ptr)); -} - -TEST_F(SlabHeapTestFixture, overallocate_2GiB) { - void *ptr{nullptr}; - size_t request_bytes{1UL << 31}; - - auto slab{slab_.get()}; - slab->malloc(&ptr, request_bytes); - - ASSERT_EQ(ptr, nullptr); -} - -TEST_F(SlabHeapTestFixture, run_all_threads_once_1_1) { - run_all_threads_once(1, 1); -} - -TEST_F(SlabHeapTestFixture, run_all_threads_once_2_1) { - run_all_threads_once(2, 1); -} - -TEST_F(SlabHeapTestFixture, run_all_threads_once_64_1) { - run_all_threads_once(64, 1); -} - -TEST_F(SlabHeapTestFixture, run_all_threads_once_128_1) { - run_all_threads_once(128, 1); -} - -TEST_F(SlabHeapTestFixture, run_all_threads_once_256_1) { - run_all_threads_once(256, 1); -} - -TEST_F(SlabHeapTestFixture, run_all_threads_once_512_1) { - run_all_threads_once(512, 1); -} - -TEST_F(SlabHeapTestFixture, run_all_threads_once_1024_1) { - run_all_threads_once(1024, 1); -} - -TEST_F(SlabHeapTestFixture, run_all_threads_once_1_2) { - run_all_threads_once(1, 2); -} - -TEST_F(SlabHeapTestFixture, run_all_threads_once_1_8) { - run_all_threads_once(1, 8); -} - -TEST_F(SlabHeapTestFixture, run_all_threads_once_1_64) { - run_all_threads_once(1, 64); -} - -TEST_F(SlabHeapTestFixture, run_all_threads_once_1_128) { - run_all_threads_once(1, 128); -} - -TEST_F(SlabHeapTestFixture, run_all_threads_once_1_256) { - run_all_threads_once(1, 256); -} - -TEST_F(SlabHeapTestFixture, run_all_threads_once_1_1024) { - run_all_threads_once(1, 1024); -} - -TEST_F(SlabHeapTestFixture, run_all_threads_once_1_2048) { - run_all_threads_once(1, 2048); -} - -TEST_F(SlabHeapTestFixture, run_all_threads_once_1_4096) { - run_all_threads_once(1, 4096); -} - -TEST_F(SlabHeapTestFixture, run_all_threads_once_1_8192) { - run_all_threads_once(1, 8192); -} - -TEST_F(SlabHeapTestFixture, run_all_threads_once_1_65536) { - run_all_threads_once(1, 65536); -} - -TEST_F(SlabHeapTestFixture, run_all_threads_once_2_2) { - run_all_threads_once(2, 2); -} - -TEST_F(SlabHeapTestFixture, run_all_threads_once_64_2) { - run_all_threads_once(64, 2); -} - -TEST_F(SlabHeapTestFixture, run_all_threads_once_256_50) { - run_all_threads_once(256, 50); -} - -TEST_F(SlabHeapTestFixture, run_all_threads_once_1024_512) { - run_all_threads_once(1024, 512); -} diff --git a/tests/unit_tests/slab_heap_gtest.hpp b/tests/unit_tests/slab_heap_gtest.hpp deleted file mode 100644 index 4c5892cb6e..0000000000 --- a/tests/unit_tests/slab_heap_gtest.hpp +++ /dev/null @@ -1,113 +0,0 @@ -/****************************************************************************** - * 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_SLAB_HEAP_GTEST_HPP -#define ROCSHMEM_SLAB_HEAP_GTEST_HPP - -#include "gtest/gtest.h" - -#include "../src/memory/slab_heap.hpp" -#include "../src/util.hpp" - -namespace rocshmem { - -/** - * @brief Datatype used by test - */ -using TYPE = uint32_t; - -/** - * @brief The bit pattern written to memory by each thread. - */ -static const TYPE THREAD_VALUE {0xAA}; - -inline __device__ -void -write_to_memory(TYPE* raw_memory) { - auto thread_idx {get_flat_block_id()}; - raw_memory[thread_idx] = THREAD_VALUE; - __threadfence(); -} - -inline __device__ -TYPE* -allocate_memory(SlabHeap* slab) { - auto block_size {get_flat_block_size()}; - - TYPE* dyn_arr {nullptr}; - size_t num_bytes {block_size * sizeof(TYPE)}; - - slab->malloc(reinterpret_cast(&dyn_arr), num_bytes); - return dyn_arr; -} - -__global__ -void -all_threads_once(SlabHeap* slab) { - auto block_mem {allocate_memory(slab)}; - write_to_memory(block_mem); -} - -class SlabHeapTestFixture : public ::testing::Test { - using SLAB_PROXY_T = SlabHeapProxy; - - public: - void - run_all_threads_once(uint32_t x_block_dim, - uint32_t x_grid_dim) { - auto slab {slab_.get()}; - - const dim3 hip_blocksize(x_block_dim, 1, 1); - const dim3 hip_gridsize(x_grid_dim, 1, 1); - - hipLaunchKernelGGL(all_threads_once, - hip_gridsize, - hip_blocksize, - 0, - nullptr, - slab); - - hipError_t return_code = hipStreamSynchronize(nullptr); - if (return_code != hipSuccess) { - printf("Failed in stream synchronize\n"); - assert(return_code == hipSuccess); - } - - TYPE* ptr {reinterpret_cast(slab->get_base_ptr())}; - - size_t number_threads {x_block_dim * x_grid_dim}; - - for (size_t i {0}; i < number_threads; i++) { - ASSERT_EQ(ptr[i], THREAD_VALUE); - } - } - - protected: - /** - * @brief Slab heap object - */ - SLAB_PROXY_T slab_ {}; -}; - -} // namespace rocshmem - -#endif // ROCSHMEM_SLAB_HEAP_GTEST_HPP diff --git a/tests/unit_tests/spin_ebo_block_mutex_gtest.cpp b/tests/unit_tests/spin_ebo_block_mutex_gtest.cpp deleted file mode 100644 index ce4bf3f511..0000000000 --- a/tests/unit_tests/spin_ebo_block_mutex_gtest.cpp +++ /dev/null @@ -1,69 +0,0 @@ -/****************************************************************************** - * 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 "spin_ebo_block_mutex_gtest.hpp" - -using namespace rocshmem; - -/***************************************************************************** - ******************************* Fixture Tests ******************************* - *****************************************************************************/ - -TEST_F(SpinEBOBlockMutexTestFixture, run_all_threads_once_1_1) { - run_all_threads_once(1, 1); -} - -TEST_F(SpinEBOBlockMutexTestFixture, run_all_threads_once_1_2) { - run_all_threads_once(1, 2); -} - -TEST_F(SpinEBOBlockMutexTestFixture, run_all_threads_once_1_8) { - run_all_threads_once(1, 8); -} - -TEST_F(SpinEBOBlockMutexTestFixture, run_all_threads_once_1_64) { - run_all_threads_once(1, 64); -} - -TEST_F(SpinEBOBlockMutexTestFixture, run_all_threads_once_1_128) { - run_all_threads_once(1, 128); -} - -TEST_F(SpinEBOBlockMutexTestFixture, run_all_threads_once_1_256) { - run_all_threads_once(1, 256); -} - -TEST_F(SpinEBOBlockMutexTestFixture, run_all_threads_once_1_1024) { - run_all_threads_once(1, 1024); -} - -TEST_F(SpinEBOBlockMutexTestFixture, run_all_threads_once_1_2048) { - run_all_threads_once(1, 2048); -} - -// TEST_F(SpinEBOBlockMutexTestFixture, run_all_threads_once_1_4096) { -// run_all_threads_once(1, 4096); -//} - -// TEST_F(SpinEBOBlockMutexTestFixture, run_all_threads_once_1_8192) { -// run_all_threads_once(1, 8192); -//} diff --git a/tests/unit_tests/spin_ebo_block_mutex_gtest.hpp b/tests/unit_tests/spin_ebo_block_mutex_gtest.hpp deleted file mode 100644 index 7d333ccad0..0000000000 --- a/tests/unit_tests/spin_ebo_block_mutex_gtest.hpp +++ /dev/null @@ -1,140 +0,0 @@ -/****************************************************************************** - * 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_SPIN_EBO_BLOCK_MUTEX_GTEST_HPP -#define ROCSHMEM_SPIN_EBO_BLOCK_MUTEX_GTEST_HPP - -#include "gtest/gtest.h" - -#include "../src/memory/hip_allocator.hpp" -#include "../src/sync/spin_ebo_block_mutex.hpp" -#include "../src/util.hpp" - -namespace rocshmem { - -inline __device__ -void -increment_counter(SpinEBOBlockMutex *mutex, - size_t *counter) { - mutex->lock(); - (*counter)++; - __threadfence(); - mutex->unlock(); -} - -__global__ -void -all_threads_once(SpinEBOBlockMutex *mutex, - size_t *counter) { - increment_counter(mutex, counter); -} - -__global__ -void -block_leader_once(SpinEBOBlockMutex *mutex, - size_t *counter) { - if (is_thread_zero_in_block()) { - increment_counter(mutex, counter); - } -} - -__global__ -void -warp_leader_once(SpinEBOBlockMutex *mutex, - size_t *counter) { - if (is_thread_zero_in_wave()) { - increment_counter(mutex, counter); - } -} - -class SpinEBOBlockMutexTestFixture : public ::testing::Test { - public: - SpinEBOBlockMutexTestFixture() { - assert(mutex_ == nullptr); - hip_allocator_.allocate((void**)&mutex_, sizeof(SpinEBOBlockMutex)); - - assert(mutex_); - new (mutex_) SpinEBOBlockMutex(true); - - assert(counter_ == nullptr); - hip_allocator_.allocate((void**)&counter_, sizeof(int)); - - assert(counter_); - *counter_ = 0; - } - - ~SpinEBOBlockMutexTestFixture() { - if (mutex_) { - hip_allocator_.deallocate(mutex_); - } - - if (counter_) { - hip_allocator_.deallocate(counter_); - } - } - - void - run_all_threads_once(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, - hip_gridsize, - hip_blocksize, - 0, - nullptr, - mutex_, - counter_); - - hipError_t return_code = hipStreamSynchronize(nullptr); - if (return_code != hipSuccess) { - printf("Failed in stream synchronize\n"); - assert(return_code == hipSuccess); - } - - size_t number_threads {x_block_dim * x_grid_dim}; - - ASSERT_EQ(*counter_, number_threads); - } - - protected: - /** - * @brief An allocator to create objects in device memory. - */ - HIPAllocator hip_allocator_ {}; - - /** - * @brief A mutex to prevent data races. - */ - SpinEBOBlockMutex *mutex_ {nullptr}; - - /** - * @brief A monotonically increasing counter to track accesses. - */ - size_t *counter_ {nullptr}; -}; - - -} // namespace rocshmem - -#endif // ROCSHMEM_SPIN_EBO_BLOCK_MUTEX_GTEST_HPP