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
[ROCm/rocshmem commit: 0fd628458c]
Cette révision appartient à :
Brandon Potter
GitHub
@@ -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_
|
||||
|
||||
@@ -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 <hip/hip_runtime.h>
|
||||
|
||||
#include <initializer_list>
|
||||
#include <ostream>
|
||||
#include <string>
|
||||
|
||||
#include "../src/device_proxy.hpp"
|
||||
#include "../src/memory/hip_allocator.hpp"
|
||||
#include "../sync/abql_block_mutex.hpp"
|
||||
|
||||
namespace rocshmem {
|
||||
|
||||
/*****************************************************************************
|
||||
******************************* FORWARD LIST ********************************
|
||||
*****************************************************************************/
|
||||
|
||||
template <typename TYPE, typename ALLOC = HIPAllocator>
|
||||
class ForwardList {
|
||||
friend class ForwardListTestFixture;
|
||||
|
||||
struct Node {
|
||||
TYPE data;
|
||||
Node* next;
|
||||
};
|
||||
|
||||
template <bool CONST>
|
||||
class Iterator;
|
||||
|
||||
public:
|
||||
using iterator = Iterator<false>;
|
||||
using const_iterator = Iterator<true>;
|
||||
|
||||
/**
|
||||
* @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 <class InputIterator>
|
||||
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<TYPE> 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<TYPE> 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 <class InputIterator>
|
||||
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<TYPE> 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 <class... Args>
|
||||
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 <class... Args>
|
||||
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 <class InputIterator>
|
||||
iterator insert_after(const_iterator position, InputIterator first,
|
||||
InputIterator last);
|
||||
|
||||
iterator insert_after(const_iterator position,
|
||||
std::initializer_list<TYPE> 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 <class Predicate>
|
||||
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 <class BinaryPredicate>
|
||||
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 <class Compare>
|
||||
void merge(ForwardList& fwdlst, Compare comp); // NOLINT
|
||||
|
||||
template <class Compare>
|
||||
void merge(ForwardList&& fwdlst, Compare comp); // NOLINT
|
||||
|
||||
/**
|
||||
* @brief Sorts elements in ForwardList, altering position in container.
|
||||
*/
|
||||
void sort();
|
||||
|
||||
template <class Compare>
|
||||
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 <typename TYPE, typename ALLOC>
|
||||
template <bool CONST>
|
||||
class ForwardList<TYPE, ALLOC>::Iterator {
|
||||
using NodeT = typename ForwardList<TYPE, ALLOC>::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 <bool Q = CONST>
|
||||
typename std::enable_if<Q, const TYPE&>::type operator*();
|
||||
|
||||
template <bool Q = CONST>
|
||||
typename std::enable_if<!Q, TYPE&>::type operator*();
|
||||
|
||||
pointer operator->();
|
||||
|
||||
template <typename ITER_TYPE>
|
||||
friend bool operator==(ITER_TYPE& a, ITER_TYPE& b); // NOLINT
|
||||
|
||||
template <typename ITER_TYPE>
|
||||
friend bool operator==(ITER_TYPE& a, ITER_TYPE b); // NOLINT
|
||||
|
||||
template <typename ITER_TYPE>
|
||||
friend bool operator!=(ITER_TYPE& a, ITER_TYPE& b); // NOLINT
|
||||
|
||||
template <typename ITER_TYPE>
|
||||
friend bool operator!=(ITER_TYPE& a, ITER_TYPE b); // NOLINT
|
||||
|
||||
private:
|
||||
NodeT* node_ptr_;
|
||||
};
|
||||
|
||||
/*****************************************************************************
|
||||
******************************* COMPARATORS *********************************
|
||||
*****************************************************************************/
|
||||
|
||||
template <class TYPE, class ALLOC>
|
||||
bool operator==(const ForwardList<TYPE, ALLOC>& lhs,
|
||||
const ForwardList<TYPE, ALLOC>& rhs);
|
||||
|
||||
template <class TYPE, class ALLOC>
|
||||
bool operator==(const ForwardList<TYPE, ALLOC>& lhs, const std::string rhs);
|
||||
|
||||
template <class TYPE, class ALLOC>
|
||||
bool operator==(const std::string lhs, const ForwardList<TYPE, ALLOC>& rhs);
|
||||
|
||||
template <class TYPE, class ALLOC>
|
||||
bool operator!=(const ForwardList<TYPE, ALLOC>& lhs,
|
||||
const ForwardList<TYPE, ALLOC>& rhs);
|
||||
|
||||
template <class TYPE, class ALLOC>
|
||||
bool operator<(const ForwardList<TYPE, ALLOC>& lhs,
|
||||
const ForwardList<TYPE, ALLOC>& rhs);
|
||||
|
||||
template <class TYPE, class ALLOC>
|
||||
bool operator<=(const ForwardList<TYPE, ALLOC>& lhs,
|
||||
const ForwardList<TYPE, ALLOC>& rhs);
|
||||
|
||||
template <class TYPE, class ALLOC>
|
||||
bool operator>(const ForwardList<TYPE, ALLOC>& lhs,
|
||||
const ForwardList<TYPE, ALLOC>& rhs);
|
||||
|
||||
template <class TYPE, class ALLOC>
|
||||
bool operator>=(const ForwardList<TYPE, ALLOC>& lhs,
|
||||
const ForwardList<TYPE, ALLOC>& rhs);
|
||||
|
||||
template <class TYPE, class ALLOC>
|
||||
void swap(ForwardList<TYPE, ALLOC>& x, ForwardList<TYPE, ALLOC>& y); // NOLINT
|
||||
|
||||
/*****************************************************************************
|
||||
******************************** STRINGIFIERS *******************************
|
||||
*****************************************************************************/
|
||||
|
||||
template <class TYPE>
|
||||
std::ostream& operator<<(std::ostream& os, ForwardList<TYPE> const& list);
|
||||
|
||||
template <class TYPE, class ALLOC>
|
||||
std::string to_string(const ForwardList<TYPE, ALLOC>& list);
|
||||
|
||||
/*****************************************************************************
|
||||
*********************************** PROXY ***********************************
|
||||
*****************************************************************************/
|
||||
|
||||
template <typename ALLOC, typename TYPE>
|
||||
class ForwardListProxy {
|
||||
using ProxyT = DeviceProxy<ALLOC, ForwardList<TYPE, ALLOC>>;
|
||||
|
||||
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<TYPE, ALLOC>* get() { return proxy_.get(); }
|
||||
|
||||
private:
|
||||
ProxyT proxy_{};
|
||||
};
|
||||
|
||||
} // namespace rocshmem
|
||||
|
||||
#endif // LIBRARY_SRC_CONTAINERS_FORWARD_LIST_HPP_
|
||||
@@ -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 <algorithm>
|
||||
#include <cassert>
|
||||
#include <string>
|
||||
|
||||
#include "forward_list.hpp"
|
||||
|
||||
namespace rocshmem {
|
||||
|
||||
/*****************************************************************************
|
||||
******************************* FORWARD LIST ********************************
|
||||
*****************************************************************************/
|
||||
|
||||
template <typename TYPE, typename ALLOC>
|
||||
ForwardList<TYPE, ALLOC>::~ForwardList() {
|
||||
while (head_ != nullptr) {
|
||||
Node* temp{head_};
|
||||
head_ = head_->next;
|
||||
allocator_.deallocate(temp);
|
||||
}
|
||||
tail_ = nullptr;
|
||||
}
|
||||
|
||||
template <typename TYPE, typename ALLOC>
|
||||
ForwardList<TYPE, ALLOC>::ForwardList(const ForwardList& fwdlst) {}
|
||||
|
||||
template <typename TYPE, typename ALLOC>
|
||||
ForwardList<TYPE, ALLOC>::ForwardList(const ForwardList& fwdlst,
|
||||
const ALLOC& alloc) {}
|
||||
|
||||
template <typename TYPE, typename ALLOC>
|
||||
ForwardList<TYPE, ALLOC>::ForwardList(std::initializer_list<TYPE> il,
|
||||
const ALLOC& allocator) {
|
||||
allocator_ = allocator;
|
||||
assign(il);
|
||||
}
|
||||
|
||||
template <typename TYPE, typename ALLOC>
|
||||
typename ForwardList<TYPE, ALLOC>::iterator
|
||||
ForwardList<TYPE, ALLOC>::begin() noexcept {
|
||||
iterator i(head_);
|
||||
return i;
|
||||
}
|
||||
|
||||
template <typename TYPE, typename ALLOC>
|
||||
typename ForwardList<TYPE, ALLOC>::const_iterator
|
||||
ForwardList<TYPE, ALLOC>::begin() const noexcept {
|
||||
const_iterator i(head_);
|
||||
return i;
|
||||
}
|
||||
|
||||
template <typename TYPE, typename ALLOC>
|
||||
typename ForwardList<TYPE, ALLOC>::iterator
|
||||
ForwardList<TYPE, ALLOC>::end() noexcept {
|
||||
iterator i(tail_);
|
||||
return i;
|
||||
}
|
||||
|
||||
template <typename TYPE, typename ALLOC>
|
||||
typename ForwardList<TYPE, ALLOC>::const_iterator
|
||||
ForwardList<TYPE, ALLOC>::end() const noexcept {
|
||||
const_iterator i(tail_);
|
||||
return i;
|
||||
}
|
||||
|
||||
template <typename TYPE, typename ALLOC>
|
||||
void ForwardList<TYPE, ALLOC>::assign(std::initializer_list<TYPE> il) {
|
||||
resize(il.size());
|
||||
std::copy_n(il.begin(), il.size(), begin());
|
||||
}
|
||||
|
||||
template <typename TYPE, typename ALLOC>
|
||||
void ForwardList<TYPE, ALLOC>::resize(size_t n) {}
|
||||
|
||||
template <typename TYPE, typename ALLOC>
|
||||
void ForwardList<TYPE, ALLOC>::resize(size_t n, const TYPE& val) {}
|
||||
|
||||
template <typename TYPE, typename ALLOC>
|
||||
void ForwardList<TYPE, ALLOC>::clear() noexcept {}
|
||||
|
||||
/*****************************************************************************
|
||||
********************************* ITERATOR **********************************
|
||||
*****************************************************************************/
|
||||
|
||||
template <typename TYPE, typename ALLOC>
|
||||
template <bool CONST>
|
||||
ForwardList<TYPE, ALLOC>::Iterator<CONST>::Iterator(NodeT* ptr)
|
||||
: node_ptr_(ptr) {}
|
||||
|
||||
template <typename TYPE, typename ALLOC>
|
||||
template <bool CONST>
|
||||
typename ForwardList<TYPE, ALLOC>::template Iterator<CONST>
|
||||
ForwardList<TYPE, ALLOC>::Iterator<CONST>::operator++() {
|
||||
if (node_ptr_) {
|
||||
node_ptr_ = node_ptr_->next;
|
||||
}
|
||||
return *this;
|
||||
}
|
||||
|
||||
template <typename TYPE, typename ALLOC>
|
||||
template <bool CONST>
|
||||
typename ForwardList<TYPE, ALLOC>::template Iterator<CONST>
|
||||
ForwardList<TYPE, ALLOC>::Iterator<CONST>::operator++(int) {
|
||||
Iterator iterator = *this;
|
||||
++*this;
|
||||
return iterator;
|
||||
}
|
||||
|
||||
template <typename TYPE, typename ALLOC>
|
||||
template <bool CONST>
|
||||
template <bool Q>
|
||||
typename std::enable_if<!Q, TYPE&>::type
|
||||
ForwardList<TYPE, ALLOC>::Iterator<CONST>::operator*() {
|
||||
return node_ptr_->data;
|
||||
}
|
||||
|
||||
template <typename TYPE, typename ALLOC>
|
||||
template <bool CONST>
|
||||
template <bool Q>
|
||||
typename std::enable_if<Q, const TYPE&>::type
|
||||
ForwardList<TYPE, ALLOC>::Iterator<CONST>::operator*() {
|
||||
return node_ptr_->data;
|
||||
}
|
||||
|
||||
template <typename TYPE, typename ALLOC>
|
||||
template <bool CONST>
|
||||
TYPE* ForwardList<TYPE, ALLOC>::Iterator<CONST>::operator->() {
|
||||
return &node_ptr_->data;
|
||||
}
|
||||
|
||||
template <typename ITER_TYPE>
|
||||
bool operator==(ITER_TYPE& a, ITER_TYPE& b) { // NOLINT
|
||||
return a.node_ptr_ == b.node_ptr_;
|
||||
}
|
||||
|
||||
template <typename ITER_TYPE>
|
||||
bool operator==(ITER_TYPE& a, ITER_TYPE b) { // NOLINT
|
||||
return a.node_ptr_ == b.node_ptr_;
|
||||
}
|
||||
|
||||
template <typename ITER_TYPE>
|
||||
bool operator!=(ITER_TYPE& a, ITER_TYPE& b) { // NOLINT
|
||||
return a.node_ptr_ != b.node_ptr_;
|
||||
}
|
||||
|
||||
template <typename ITER_TYPE>
|
||||
bool operator!=(ITER_TYPE& a, ITER_TYPE b) { // NOLINT
|
||||
return a.node_ptr_ != b.node_ptr_;
|
||||
}
|
||||
|
||||
/*****************************************************************************
|
||||
******************************* COMPARATORS *********************************
|
||||
*****************************************************************************/
|
||||
|
||||
template <typename TYPE, typename ALLOC>
|
||||
bool operator==(const ForwardList<TYPE, ALLOC>& lhs,
|
||||
const ForwardList<TYPE, ALLOC>& rhs) {
|
||||
return false;
|
||||
}
|
||||
|
||||
template <typename TYPE, typename ALLOC>
|
||||
bool operator==(const ForwardList<TYPE, ALLOC>& lhs, const std::string rhs) {
|
||||
return false;
|
||||
}
|
||||
|
||||
template <typename TYPE, typename ALLOC>
|
||||
bool operator==(const std::string lhs, const ForwardList<TYPE, ALLOC>& rhs) {
|
||||
return false;
|
||||
}
|
||||
|
||||
template <typename TYPE, typename ALLOC>
|
||||
bool operator!=(const ForwardList<TYPE, ALLOC>& lhs,
|
||||
const ForwardList<TYPE, ALLOC>& rhs) {
|
||||
return false;
|
||||
}
|
||||
|
||||
template <typename TYPE, typename ALLOC>
|
||||
bool operator!=(const ForwardList<TYPE, ALLOC>& lhs, const std::string rhs) {
|
||||
return false;
|
||||
}
|
||||
|
||||
template <typename TYPE, typename ALLOC>
|
||||
bool operator!=(const std::string lhs, const ForwardList<TYPE, ALLOC>& rhs) {
|
||||
return false;
|
||||
}
|
||||
|
||||
/*****************************************************************************
|
||||
******************************** STRINGIFIERS *******************************
|
||||
*****************************************************************************/
|
||||
|
||||
template <typename TYPE>
|
||||
std::ostream& operator<<(std::ostream& os, ForwardList<TYPE> const& list) {
|
||||
using CItr = typename ForwardList<TYPE>::const_iterator;
|
||||
for (CItr ci = list.begin(); ci != list.end(); ++ci) {
|
||||
os << *ci << " ";
|
||||
}
|
||||
return os;
|
||||
}
|
||||
|
||||
template <typename TYPE, typename ALLOC>
|
||||
std::string to_string(const ForwardList<TYPE, ALLOC>& list) {
|
||||
std::stringstream ss;
|
||||
ss << list;
|
||||
return ss.str();
|
||||
}
|
||||
|
||||
} // namespace rocshmem
|
||||
|
||||
#endif // LIBRARY_SRC_CONTAINERS_FORWARD_LIST_IMPL_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 {
|
||||
|
||||
@@ -27,6 +27,5 @@ target_sources(
|
||||
${PROJECT_NAME}
|
||||
PRIVATE
|
||||
single_heap.cpp
|
||||
slab_heap.cpp
|
||||
memory_allocator.cpp
|
||||
)
|
||||
|
||||
@@ -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 <sstream>
|
||||
|
||||
#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<char**>(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<char**>(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<uint64_t>(*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<char*>(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<char*>(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
|
||||
@@ -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<HEAP_T>;
|
||||
|
||||
/**
|
||||
* @brief Helper type for notifier
|
||||
*/
|
||||
using NOTIFIER_PROXY_T = NotifierProxy<HIPAllocator, detail::atomic::memory_scope_workgroup>;
|
||||
|
||||
/**
|
||||
* @brief Helper type for notifier
|
||||
*/
|
||||
using MUTEX_PROXY_T = ABQLBlockMutexProxy<HIPAllocator>;
|
||||
|
||||
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 <typename ALLOCATOR>
|
||||
class SlabHeapProxy {
|
||||
using ProxyT = DeviceProxy<ALLOCATOR, SlabHeap>;
|
||||
|
||||
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_
|
||||
@@ -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};
|
||||
|
||||
@@ -26,6 +26,5 @@
|
||||
target_sources(
|
||||
${PROJECT_NAME}
|
||||
PRIVATE
|
||||
spin_ebo_block_mutex.cpp
|
||||
abql_block_mutex.cpp
|
||||
)
|
||||
|
||||
@@ -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 <algorithm>
|
||||
|
||||
#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<int *>(&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<int *>(&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<int *>(&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
|
||||
@@ -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 <hip/hip_runtime.h>
|
||||
|
||||
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_
|
||||
@@ -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) {
|
||||
|
||||
@@ -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;
|
||||
|
||||
@@ -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
|
||||
)
|
||||
|
||||
@@ -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
|
||||
*/
|
||||
}
|
||||
@@ -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
|
||||
@@ -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 <iostream>
|
||||
#include <string>
|
||||
|
||||
using namespace rocshmem;
|
||||
|
||||
/*****************************************************************************
|
||||
******************************* Fixture Tests *******************************
|
||||
*****************************************************************************/
|
||||
|
||||
TEST_F(ForwardListTestFixture, default_constructor) {
|
||||
default_constructor_test();
|
||||
}
|
||||
|
||||
TEST(ForwardListTest, constructor_tests) {
|
||||
ForwardList<std::string> list_1{"rocshmem", "forward_list"};
|
||||
std::string str_1 = "rocshmem forward_list";
|
||||
ASSERT_EQ(to_string(list_1), str_1);
|
||||
|
||||
// ForwardList<std::string> list_2(list_1.begin(), list_1.end());
|
||||
// std::cout << "list_2: " << list_2 << '\n';
|
||||
|
||||
// ForwardList<std::string> list_3(list_1);
|
||||
// std::cout << "list_3: " << list_3 << '\n';
|
||||
|
||||
// ForwardList<std::string> list_4(5, "rocm");
|
||||
// std::cout << "list_4: " << list_4 << '\n';
|
||||
}
|
||||
@@ -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<std::string> list_ {};
|
||||
};
|
||||
|
||||
} // namespace rocshmem
|
||||
|
||||
#endif // ROCSHMEM_FORWARD_LIST_GTEST_HPP
|
||||
@@ -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);
|
||||
}
|
||||
@@ -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<void**>(&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<HIPAllocator>;
|
||||
|
||||
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<TYPE*>(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
|
||||
@@ -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);
|
||||
//}
|
||||
@@ -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
|
||||
Référencer dans un nouveau ticket
Bloquer un utilisateur