diff --git a/cmake/SetupChaiOptions.cmake b/cmake/SetupChaiOptions.cmake index d7b69ac6..3ead3f31 100644 --- a/cmake/SetupChaiOptions.cmake +++ b/cmake/SetupChaiOptions.cmake @@ -20,6 +20,7 @@ option(CHAI_ENABLE_MANAGED_PTR "Enable managed_ptr" On) option(CHAI_DEBUG "Enable Debug Logging." Off) option(CHAI_ENABLE_RAJA_NESTED_TEST "Enable raja-chai-nested-tests, which fails to build on Debug CUDA builds." On) option(CHAI_ENABLE_MANAGED_PTR_ON_GPU "Enable managed_ptr on GPU" On) +option(CHAI_ENABLE_EXPERIMENTAL "Enable experimental features" On) option(CHAI_ENABLE_TESTS "Enable CHAI tests" On) option(CHAI_ENABLE_BENCHMARKS "Enable benchmarks" Off) diff --git a/host-configs/lc/blueos_3_ppc64le_ib_p9/nvcc_clang.cmake b/host-configs/lc/blueos_3_ppc64le_ib_p9/nvcc_clang.cmake index 4b2ae845..fdd69246 100644 --- a/host-configs/lc/blueos_3_ppc64le_ib_p9/nvcc_clang.cmake +++ b/host-configs/lc/blueos_3_ppc64le_ib_p9/nvcc_clang.cmake @@ -30,3 +30,5 @@ set(CMAKE_CUDA_COMPILER "${CUDA_TOOLKIT_ROOT_DIR}/bin/nvcc" CACHE PATH "") set(CMAKE_CUDA_HOST_COMPILER "${CMAKE_CXX_COMPILER}" CACHE PATH "") set(CMAKE_CUDA_ARCHITECTURES "70" CACHE STRING "") set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -Xcompiler=--gcc-toolchain=${GCC_HOME}" CACHE STRING "") + +set(UMPIRE_FMT_TARGET "fmt::fmt" CACHE STRING "") diff --git a/host-configs/lc/toss_4_x86_64_ib_cray/amdclang.cmake b/host-configs/lc/toss_4_x86_64_ib_cray/amdclang.cmake index db1dafbd..ae0e9bc1 100644 --- a/host-configs/lc/toss_4_x86_64_ib_cray/amdclang.cmake +++ b/host-configs/lc/toss_4_x86_64_ib_cray/amdclang.cmake @@ -6,8 +6,8 @@ ############################################################################## # Set up software versions -set(ROCM_VERSION "6.2.0" CACHE PATH "") -set(GCC_VERSION "12.2.1" CACHE PATH "") +set(ROCM_VERSION "6.4.1" CACHE PATH "") +set(GCC_VERSION "13.3.1" CACHE PATH "") # Set up compilers set(COMPILER_BASE "/usr/tce/packages/rocmcc/rocmcc-${ROCM_VERSION}-magic" CACHE PATH "") diff --git a/src/chai/CMakeLists.txt b/src/chai/CMakeLists.txt index faab5481..c8443eeb 100644 --- a/src/chai/CMakeLists.txt +++ b/src/chai/CMakeLists.txt @@ -22,6 +22,13 @@ set (chai_headers PointerRecord.hpp Types.hpp) +if(CHAI_ENABLE_EXPERIMENTAL) + set(chai_headers + ${chai_headers} + expt/ExecutionContext.hpp + expt/ExecutionContextManager.hpp) +endif() + if(CHAI_DISABLE_RM) set(chai_headers ${chai_headers} diff --git a/src/chai/expt/Allocator.hpp b/src/chai/expt/Allocator.hpp new file mode 100644 index 00000000..e180cd47 --- /dev/null +++ b/src/chai/expt/Allocator.hpp @@ -0,0 +1,77 @@ +#ifndef CHAI_ALLOCATOR_HPP +#define CHAI_ALLOCATOR_HPP + +namespace chai::expt { + class Allocator { + private: + class AllocatorConcept + { + public: + virtual ~AllocatorConcept() = default; + virtual void* do_allocate(std::size_t bytes) = 0; + virtual void do_deallocate(void* ptr) = 0; + virtual std::unique_ptr clone() const = 0; + }; // class AllocatorConcept + + template + class AllocatorModel : public AllocatorConcept + { + public: + AllocatorModel(AllocatorType allocator) + : m_allocator{std::move(allocator)} + { + } + + virtual void* allocate(std::size_t bytes) override + { + return allocate(m_allocator, bytes); + } + + virtual void do_deallocate(void* ptr) override + { + deallocate(m_allocator, ptr); + } + + virtual std::unique_ptr clone() const override + { + return std::make_unique(*this); + } + + private: + AllocatorType m_allocator; + }; // class AllocatorModel + + friend void* allocate(const Allocator& allocator, std::size_t bytes) + { + return allocator.m_pimpl->do_allocate(bytes); + } + + friend void deallocate(const Allocator& allocator, void* ptr) + { + allocator.m_pimple->do_deallocate(ptr); + } + + std::unique_ptr m_pimpl; + + public: + template + Allocator(AllocatorType allocator) + : m_pimpl{std::make_unique>(std::move(allocator))} + { + } + + Allocator(const Allocator& other) + : m_pimple{other.m_pimpl->clone()} + { + } + + Allocator& operator=(const Allocator& other) + { + Allocator temp(other); + std::swap(m_pimpl, temp.m_pimpl); + return *this; + } + }; // class Allocator +} + +#endif // CHAI_ALLOCATOR_HPP \ No newline at end of file diff --git a/src/chai/expt/Array.hpp b/src/chai/expt/Array.hpp new file mode 100644 index 00000000..db452a37 --- /dev/null +++ b/src/chai/expt/Array.hpp @@ -0,0 +1,275 @@ +#ifndef CHAI_MANAGED_ARRAY_HPP +#define CHAI_MANAGED_ARRAY_HPP + +#include "chai/expt/ArrayManager.hpp" +#include "chai/expt/ExecutionContextManager.hpp" +#include + +namespace chai { +namespace expt { + /*! + * \class ManagedArray + * + * \brief An array class that manages coherency across the CPU and GPU. + * How the coherence is obtained is controlled by the array manager. + * + * \tparam ElementType The type of element in the array. + */ + template + class ManagedArray { + public: + /*! + * \brief Constructs an empty array without an array manager. + */ + ManagedArray() = default; + + /*! + * \brief Constructs an array from a manager. + * + * \param manager The array manager controls the coherence of the array. + * + * \note The array takes ownership of the manager. + */ + explicit ManagedArray(ArrayManager* manager) + : m_manager{manager} + { + if (m_manager) + { + m_size = m_manager->size(); + } + } + + /*! + * \brief Constructs a shallow copy of an array from another and makes + * the data coherent in the current execution space. + * + * \param other The other array. + * + * \note This is a shallow copy. + */ + CHAI_HOST_DEVICE ManagedArray(const ManagedArray& other) : + m_data{other.m_data}, + m_size{other.m_size}, + m_manager{other.m_manager} + { +#if !defined(CHAI_DEVICE_COMPILE) + if (m_manager) { + m_data = m_manager->data(!std::is_const::value)); + } +#endif + } + + /*! + * \brief Constructs a ManagedArray from a ManagedArray. + * + * \param other The non-const array to convert from. + * + * \note This is a converting constructor that enables implicit conversion + * from ManagedArray to ManagedArray. + */ + template * = nullptr> + CHAI_HOST_DEVICE ManagedArray(const ManagedArray& other) + : m_data{other.m_data}, + m_size{other.m_size}, + m_manager{other.m_manager} + { + } + + /*! + * \brief Sets the array manager for this ManagedArray. + * + * \param manager The new array manager to be set. + * + * \post The ManagedArray takes ownership of the new manager objet. + */ + void setManager(ArrayManager* manager) + { + delete m_manager; + m_manager = manager; + } + + /*! + * \brief Get the array manager associated with this ManagedArray. + * + * \return A pointer to the array manager. + */ + ArrayManager* getManager() const { + return m_manager; + } + + /*! + * \brief Resizes the array to the specified new size. + * + * \param newSize The new size to resize the array to. + * + * \note This method updates the size of the array and triggers a resize operation in the array manager if it exists. + * If no array manager is associated, an exception is thrown. + */ + void resize(std::size_t newSize) { + if (m_manager) { + m_size = newSize; + m_manager->resize(newSize); + } + else { + throw std::runtime_exception("Unable to resize"); + } + } + + /*! + * \brief Frees the resources associated with this array. + * + * \note Once free has been called, it is invalid to use any other copies + * of this array (since copies are shallow). + */ + void free() { + m_data = nullptr; + m_size = 0; + delete m_manager; + m_manager = nullptr; + } + + /*! + * \brief Get the number of elements in the array. + * + * \pre The copy constructor has been called with the execution space + * set to CPU or GPU (e.g. by the RAJA plugin). + */ + CHAI_HOST_DEVICE std::size_t size() const { + return m_size; + } + + /*! + * \brief Get a pointer to the element data in the specified context. + * + * \param context The context in which to retrieve the element data. + * + * \return A pointer to the element data in the specified context. + */ + ElementType* data(ExecutionContext context) const { + if (m_manager) { + m_data = m_manager->data(context, !std::is_const::value); + } + + return m_data; + } + + /*! + * \brief Get a const pointer to the element data in the specified context. + * + * \param context The context in which to retrieve the const element data. + * + * \return A const pointer to the element data in the specified context. + */ + const ElementType* cdata(ExecutionContext context) const { + if (m_manager) { + m_data = m_manager->data(context, false); + } + + return m_data; + } + + /*! + * \brief Get a pointer to the element data in the current execution space. + * + * \return A pointer to the element data in the current execution space. + */ + CHAI_HOST_DEVICE ElementType* data() const { +#if !defined(CHAI_DEVICE_COMPILE) + return data(HOST); +#endif + return m_data; + } + + /*! + * \brief Get a const pointer to the element data in the current execution space. + * + * \return A const pointer to the element data in the current execution space. + */ + CHAI_HOST_DEVICE const ElementType* cdata() const { +#if !defined(CHAI_DEVICE_COMPILE) + return cdata(HOST); +#endif + return m_data; + } + + /*! + * \brief Get the ith element in the array. + * + * \param i The index of the element to retrieve. + * + * \pre The copy constructor has been called with the execution space + * set to CPU or GPU (e.g. by the RAJA plugin). + */ + CHAI_HOST_DEVICE ElementType& operator[](std::size_t i) const { + return m_data[i]; + } + + /*! + * \brief Get the value of the element at the specified index. + * + * \param i The index of the element to retrieve. + * + * \return The value of the element at the specified index. + * + * \throw std::runtime_exception if unable to retrieve the element. + */ + ElementType get(std::size_t i) const { + if (m_manager) { + return m_manager->get(i); + } + else { + throw std::runtime_exception("Unable to get element"); + } + } + + /*! + * \brief Set a value at a specified index in the array. + * + * \param i The index where the value is to be set. + * \param value The value to set at the specified index. + * + * \throw std::runtime_exception if the array manager is not associated with the ManagedArray. + */ + void set(std::size_t i, const ElementType& value) { + if (m_manager) { + m_manager->set(i, value); + } + else { + throw std::runtime_exception("Unable to set element"); + } + } + + private: + /*! + * The array that is coherent in the current execution space. + */ + ElementType* m_data = nullptr; + + /*! + * The number of elements in the array. + */ + std::size_t m_size = 0; + + /*! + * The array manager controls the coherence of the array. + */ + ArrayManager* m_manager = nullptr; + }; // class ManagedArray + + /*! + * \brief Constructs an array by creating a new manager object. + * + * \tparam ArrayManager The type of array manager. + * \tparam Args The type of the arguments used to construct the array manager. + * + * \param args The arguments to construct an array manager. + */ + template , typename... Args> + ManagedArray makeArray(Args&&... args) { + return ManagedArray(new ArrayManager(std::forward(args)...)); + } +} // namespace expt +} // namespace chai + +#endif // CHAI_MANAGED_ARRAY_HPP diff --git a/src/chai/expt/ArrayManager.hpp b/src/chai/expt/ArrayManager.hpp new file mode 100644 index 00000000..e98e92bd --- /dev/null +++ b/src/chai/expt/ArrayManager.hpp @@ -0,0 +1,79 @@ +////////////////////////////////////////////////////////////////////////////// +// Copyright (c) 2016-25, Lawrence Livermore National Security, LLC and CHAI +// project contributors. See the CHAI LICENSE file for details. +// +// SPDX-License-Identifier: BSD-3-Clause +////////////////////////////////////////////////////////////////////////////// + +#ifndef CHAI_ARRAY_MANAGER_HPP +#define CHAI_ARRAY_MANAGER_HPP + +#include "chai/expt/Context.hpp" +#include + +namespace chai { +namespace expt { + /*! + * \class ArrayManager + * + * \brief Controls the coherence of an array. + */ + class ArrayManager { + public: + /*! + * \brief Virtual destructor. + */ + virtual ~ArrayManager() = default; + + /*! + * \brief Creates a clone of this ArrayManager. + * + * \return A new ArrayManager object that is a clone of this instance. + */ + virtual ArrayManager* clone() const = 0; + + /*! + * \brief Resizes the array to the specified new size. + * + * \param newSize The new size to resize the array to. + */ + virtual void resize(std::size_t newSize) = 0; + + /*! + * \brief Returns the size of the contained array. + * + * \return The size of the contained array. + */ + virtual std::size_t size() const = 0; + + /*! + * \brief Updates the data to be coherent in the current execution context. + * + * \param data [out] A coherent array in the current execution context. + */ + virtual void* data(bool touch) = 0; + + /*! + * \brief Returns the value at index i. + * + * Note: Use this function sparingly as it may be slow. + * + * \param i The index of the element to get. + * \return The value at index i. + */ + virtual void* get(std::size_t offset, std::size_t size) const = 0; + + /*! + * \brief Sets the value at index i to the specified value. + * + * Note: Use this function sparingly as it may be slow. + * + * \param i The index of the element to set. + * \param value The value to set at index i. + */ + virtual void set(std::size_t offset, std::size_t size, const void* value) = 0; + }; // class ArrayManager +} // namespace expt +} // namespace chai + +#endif // CHAI_ARRAY_MANAGER_HPP diff --git a/src/chai/expt/CopyHidingArray.hpp b/src/chai/expt/CopyHidingArray.hpp new file mode 100644 index 00000000..b22e27be --- /dev/null +++ b/src/chai/expt/CopyHidingArray.hpp @@ -0,0 +1,418 @@ +#ifndef CHAI_COPY_HIDING_ARRAY_HPP +#define CHAI_COPY_HIDING_ARRAY_HPP + +#include "umpire/ResourceManager.hpp" + +// TODO: Determine how to specify starting execution space + +namespace chai { +namespace expt { + /*! + * \class CopyHidingArray + * + * \brief Controls the coherence of an array on the host and device. + */ + template + class CopyHidingArray + public: + /*! + * Constructs a CopyHidingArray with default allocators from Umpire + * for the "HOST" and "DEVICE" resources. + */ + CopyHidingArray() = default; + + /*! + * Constructs a CopyHidingArray with the given Umpire allocators. + */ + CopyHidingArray(const umpire::Allocator& cpuAllocator, + const umpire::Allocator& gpuAllocator) : + m_cpu_allocator{cpuAllocator}, + m_gpu_allocator{gpuAllocator} + { + } + + /*! + * Constructs a CopyHidingArray with the given Umpire allocator IDs. + */ + CopyHidingArray(int cpuAllocatorID, + int gpuAllocatorID) : + m_resource_manager{umpire::ResourceManager::getInstance()}, + m_cpu_allocator{m_resource_manager.getAllocator(cpuAllocatorID)}, + m_gpu_allocator{m_resource_manager.getAllocator(gpuAllocatorID)} + { + } + + /*! + * Constructs a CopyHidingArray with the given size using default allocators + * from Umpire for the "HOST" and "DEVICE" resources. + */ + CopyHidingArray(size_type size) : + m_size{size} + { + // TODO: Exception handling + m_cpu_data = m_cpu_allocator.allocate(size); + m_gpu_data = m_gpu_allocator.allocate(size); + } + + /*! + * Constructs a CopyHidingArray with the given size using the given Umpire + * allocators. + */ + CopyHidingArray(size_type size, + const umpire::Allocator& cpuAllocator, + const umpire::Allocator& gpuAllocator) : + m_cpu_allocator{cpuAllocator}, + m_gpu_allocator{gpuAllocator}, + m_size{size} + { + // TODO: Exception handling + m_cpu_data = m_cpu_allocator.allocate(size); + m_gpu_data = m_gpu_allocator.allocate(size); + } + + /*! + * Constructs a CopyHidingArray with the given size using the given Umpire + * allocator IDs. + */ + CopyHidingArray(size_type size, + int cpuAllocatorID, + int gpuAllocatorID) : + m_resource_manager{umpire::ResourceManager::getInstance()}, + m_cpu_allocator{m_resource_manager.getAllocator(cpuAllocatorID)}, + m_gpu_allocator{m_resource_manager.getAllocator(gpuAllocatorID)}, + m_size{size} + { + // TODO: Exception handling + m_cpu_data = m_cpu_allocator.allocate(size); + m_gpu_data = m_gpu_allocator.allocate(size); + } + + /*! + * Constructs a deep copy of the given CopyHidingArray. + */ + CopyHidingArray(const CopyHidingArray& other) : + m_cpu_allocator{other.m_cpu_allocator}, + m_gpu_allocator{other.m_gpu_allocator}, + m_size{other.m_size}, + m_touch{other.m_touch} + { + if (other.m_cpu_data) + { + m_cpu_data = m_cpu_allocator.allocate(m_size); + m_resourceManager.copy(m_cpu_data, other.m_cpu_data, m_size); + } + + if (other.m_gpu_data) + { + m_gpu_data = m_gpu_allocator.allocate(m_size); + m_resourceManager.copy(m_gpu_data, other.m_gpu_data, m_size); + } + } + + /*! + * Constructs a CopyHidingArray that takes ownership of the + * resources from the given CopyHidingArray. + */ + CopyHidingArray(CopyHidingArray&& other) : + m_cpu_allocator{other.m_cpu_allocator}, + m_gpu_allocator{other.m_gpu_allocator}, + m_size{other.m_size}, + m_touch{other.m_touch}, + m_cpu_data{other.m_cpu_data}, + m_gpu_data{other.m_gpu_data} + { + other.m_size = 0; + other.m_cpu_data = nullptr; + other.m_gpu_data = nullptr; + other.m_touch = ExecutionContext::NONE; + } + + /*! + * \brief Virtual destructor. + */ + ~CopyHidingArray() + { + m_cpu_allocator.deallocate(m_cpu_data); + m_gpu_allocator.deallocate(m_gpu_data); + } + + /*! + * \brief Copy assignment operator. + */ + CopyHidingArray& operator=(const CopyHidingArray& other) + { + if (this != &other) + { + // Copy-assign base class if needed (uncomment if Manager is copy-assignable) + // Manager::operator=(other); + + // Copy-assign or copy members + m_cpu_allocator = other.m_cpu_allocator; + m_gpu_allocator = other.m_gpu_allocator; + m_touch = other.m_touch; + + // Allocate new resources before releasing old ones for strong exception safety + void* new_cpu_data = nullptr; + void* new_gpu_data = nullptr; + + if (other.m_cpu_data) + { + new_cpu_data = m_cpu_allocator.allocate(other.m_size); + m_resourceManager.copy(new_cpu_data, other.m_cpu_data, other.m_size); + } + + if (other.m_gpu_data) + { + new_gpu_data = m_gpu_allocator.allocate(other.m_size); + m_resourceManager.copy(new_gpu_data, other.m_gpu_data, other.m_size); + } + + // Clean up old resources + if (m_cpu_data) + { + m_cpu_allocator.deallocate(m_cpu_data, m_size); + } + + if (m_gpu_data) + { + m_gpu_allocator.deallocate(m_gpu_data, m_size); + } + + // Assign new resources and size + m_cpu_data = new_cpu_data; + m_gpu_data = new_gpu_data; + m_size = other.m_size; + } + + return *this; + } + + /*! + * \brief Move assignment operator. + */ + CopyHidingArray& operator=(CopyHidingArray&& other) + { + if (this != &other) + { + // Release any resources currently held + if (m_cpu_data) + { + m_cpu_allocator.deallocate(m_cpu_data, m_size); + m_cpu_data = nullptr; + } + + if (m_gpu_data) + { + m_gpu_allocator.deallocate(m_gpu_data, m_size); + m_gpu_data = nullptr; + } + + // Move-assign base class if needed (uncomment if Manager is move-assignable) + // Manager::operator=(std::move(other)); + + // Move-assign or copy members + m_cpu_allocator = other.m_cpu_allocator; + m_gpu_allocator = other.m_gpu_allocator; + m_size = other.m_size; + m_cpu_data = other.m_cpu_data; + m_gpu_data = other.m_gpu_data; + m_touch = other.m_touch; + + // Null out other's pointers and reset size + other.m_cpu_data = nullptr; + other.m_gpu_data = nullptr; + other.m_size = 0; + other.m_touch = ExecutionContext::NONE; + } + + return *this; + } + + /*! + * \brief Resize the underlying arrays. + */ + void resize(size_type newSize) + { + if (newSize != m_size) + { + if (m_touch == ExecutionContext::CPU) + { + m_resource_manager.reallocate(m_cpu_pointer, newSize); + + if (m_gpu_pointer) + { + m_resource_manager.deallocate(m_gpu_pointer); + m_gpu_pointer = m_gpu_allocator.allocate(newSize); + } + } + else if (m_touch == ExecutionContext::GPU) + { + m_resource_manager.reallocate(m_gpu_pointer, newSize); + + if (m_cpu_pointer) + { + m_resource_manager.deallocate(m_cpu_pointer); + m_cpu_pointer = m_cpu_allocator.allocate(newSize); + } + } + else + { + if (m_gpu_pointer) + { + m_resource_manager.reallocate(m_gpu_pointer, newSize); + } + + if (m_cpu_pointer) + { + m_resource_manager.reallocate(m_cpu_pointer, newSize); + } + } + } + } + + /*! + * \brief Get the size of the underlying arrays. + */ + size_type size() const + { + return m_size; + } + + /*! + * \brief Updates the data to be coherent in the current execution space. + */ + T* data(ExecutionContext context) + { + if (context == ExecutionContext::CPU) + { + if (!m_cpu_data) + { + m_cpu_data = m_cpu_allocator.allocate(m_size); + } + + if (m_touch == ExecutionContext::GPU) + { + m_resource_manager.copy(m_cpu_data, m_gpu_data, m_size); + m_touch = ExecutionContext::NONE; + } + + if (touch) + { + m_touch = ExecutionContext::CPU; + } + + return m_cpu_data; + } + else if (context == ExecutionContext::GPU) + { + if (!m_gpu_data) + { + m_gpu_data = m_gpu_allocator.allocate(m_size); + } + + if (m_touch == ExecutionContext::CPU) + { + m_resource_manager.copy(m_gpu_data, m_cpu_data, m_size); + m_touch = ExecutionContext::NONE; + } + + if (touch) + { + m_touch = ExecutionContext::GPU; + } + + return m_gpu_data; + } + else + { + return nullptr; + } + } + + /*! + * \brief Updates the data to be coherent in the current execution space. + */ + const T* data(ExecutionContext context) const + { + if (context == ExecutionContext::CPU) + { + if (!m_cpu_data) + { + m_cpu_data = m_cpu_allocator.allocate(m_size); + } + + if (m_touch == ExecutionContext::GPU) + { + m_resource_manager.copy(m_cpu_data, m_gpu_data, m_size); + m_touch = ExecutionContext::NONE; + } + + return m_cpu_data; + } + else if (context == ExecutionContext::GPU) + { + if (!m_gpu_data) + { + m_gpu_data = m_gpu_allocator.allocate(m_size); + } + + if (m_touch == ExecutionContext::CPU) + { + m_resource_manager.copy(m_gpu_data, m_cpu_data, m_size); + m_touch = ExecutionContext::NONE; + } + + return m_gpu_data; + } + else + { + return nullptr; + } + } + +#if 0 + /*! + * \brief Get the i-th element. + * + * \warning Use sparingly, as coherence must be checked. + */ + ElementType getElement(size_type i) const + { + if (m_touch == ExecutionContext::GPU) + { + // Copy m_gpu_data[i] to host + } + else + { + return m_cpu_data[i]; + } + + return m_cpu_data[i]; + } + + void setElement(size_type i, const ElementType& value) + { + if (m_touch == ExecutionContext::GPU) + { + // Copy value to m_gpu_data[i] + } + else + { + m_cpu_data[i] = value; + } + } +#endif + + private: + umpire::ResourceManager& m_resource_manager{umpire::ResourceManager::getInstance()}; + umpire::Allocator m_cpu_allocator{m_resource_manager.getAllocator("HOST")}; + umpire::Allocator m_gpu_allocator{m_resource_manager.getAllocator("DEVICE")}; + size_type m_size{0}; + ElementType* m_cpu_data{nullptr}; + ElementType* m_gpu_data{nullptr}; + ExecutionContext m_touch{ExecutionContext::NONE}; + }; // class CopyHidingArray +} // namespace expt +} // namespace chai + +#endif // CHAI_COPY_HIDING_ARRAY_HPP diff --git a/src/chai/expt/CopyHidingArrayManager.hpp b/src/chai/expt/CopyHidingArrayManager.hpp new file mode 100644 index 00000000..802fe90a --- /dev/null +++ b/src/chai/expt/CopyHidingArrayManager.hpp @@ -0,0 +1,411 @@ +////////////////////////////////////////////////////////////////////////////// +// Copyright (c) 2016-25, Lawrence Livermore National Security, LLC and CHAI +// project contributors. See the CHAI LICENSE file for details. +// +// SPDX-License-Identifier: BSD-3-Clause +////////////////////////////////////////////////////////////////////////////// + +#ifndef CHAI_COPY_HIDING_ARRAY_MANAGER_HPP +#define CHAI_COPY_HIDING_ARRAY_MANAGER_HPP + +#include "chai/expt/ArrayManager.hpp" +#include "chai/expt/ContextManager.hpp" +#include "umpire/ResourceManager.hpp" + +namespace chai { +namespace expt { + /*! + * \class CopyHidingArrayManager + * + * \brief Controls the coherence of an array on the host and device. + */ + template + class CopyHidingArrayManager : public ArrayManager { + public: + /*! + * Constructs a CopyHidingArrayManager with default allocators from Umpire + * for the "HOST" and "DEVICE" resources. + */ + CopyHidingArrayManager() = default; + + /*! + * Constructs a CopyHidingArrayManager with the given Umpire allocators. + */ + CopyHidingArrayManager(const umpire::Allocator& hostAllocator, + const umpire::Allocator& deviceAllocator) + : ArrayManager{}, + m_host_allocator{hostAllocator}, + m_device_allocator{deviceAllocator} + { + } + + /*! + * Constructs a CopyHidingArrayManager with the given Umpire allocator IDs. + */ + CopyHidingArrayManager(int hostAllocatorID, + int deviceAllocatorID) + : ArrayManager{}, + m_resource_manager{umpire::ResourceManager::getInstance()}, + m_host_allocator{m_resource_manager.getAllocator(hostAllocatorID)}, + m_device_allocator{m_resource_manager.getAllocator(deviceAllocatorID)} + { + } + + /*! + * Constructs a CopyHidingArrayManager with the given size using default allocators + * from Umpire for the "HOST" and "DEVICE" resources. + */ + CopyHidingArrayManager(std::size_t size) + : ArrayManager{}, + m_size{size} + { + // TODO: Exception handling + m_host_data = static_cast(m_host_allocator.allocate(size*sizeof(ElementT)); + m_device_data = static_cast(m_device_allocator.allocate(size*sizeof(ElementT)); + } + + /*! + * Constructs a CopyHidingArrayManager with the given size using the given Umpire + * allocators. + */ + CopyHidingArrayManager(std::size_t size, + const umpire::Allocator& hostAllocator, + const umpire::Allocator& deviceAllocator) + : ArrayManager{}, + m_host_allocator{hostAllocator}, + m_device_allocator{deviceAllocator}, + m_size{size} + { + // TODO: Exception handling + m_host_data = m_host_allocator.allocate(size); + m_device_data = m_device_allocator.allocate(size); + } + + /*! + * Constructs a CopyHidingArrayManager with the given size using the given Umpire + * allocator IDs. + */ + CopyHidingArrayManager(std::size_t size, + int hostAllocatorID, + int deviceAllocatorID) + : ArrayManager{}, + m_resource_manager{umpire::ResourceManager::getInstance()}, + m_host_allocator{m_resource_manager.getAllocator(hostAllocatorID)}, + m_device_allocator{m_resource_manager.getAllocator(deviceAllocatorID)}, + m_size{size} + { + // TODO: Exception handling + m_host_data = m_host_allocator.allocate(size); + m_device_data = m_device_allocator.allocate(size); + } + + /*! + * Constructs a deep copy of the given CopyHidingArrayManager. + */ + CopyHidingArrayManager(const CopyHidingArrayManager& other) + : ArrayManager{}, + m_host_allocator{other.m_host_allocator}, + m_device_allocator{other.m_device_allocator}, + m_size{other.m_size}, + m_touch{other.m_touch} + { + if (other.m_host_data) + { + m_host_data = m_host_allocator.allocate(m_size); + m_resource_manager.copy(m_host_data, other.m_host_data, m_size*sizeof(ElementT)); + } + + if (other.m_device_data) + { + m_device_data = m_device_allocator.allocate(m_size); + m_resource_manager.copy(m_device_data, other.m_device_data, m_size*sizeof(ElementT)); + } + } + + /*! + * Constructs a CopyHidingArrayManager that takes ownership of the + * resources from the given CopyHidingArrayManager. + */ + CopyHidingArrayManager(CopyHidingArrayManager&& other) noexcept + : ArrayManager{}, + m_host_allocator{other.m_host_allocator}, + m_device_allocator{other.m_device_allocator}, + m_size{other.m_size}, + m_touch{other.m_touch}, + m_host_data{other.m_host_data}, + m_device_data{other.m_device_data} + { + other.m_size = 0; + other.m_host_data = nullptr; + other.m_device_data = nullptr; + other.m_touch = NONE; + } + + /*! + * \brief Virtual destructor. + */ + virtual ~CopyHidingArrayManager() + { + if (m_host_data) { + m_host_allocator.deallocate(m_host_data); + } + if (m_device_data) { + m_device_allocator.deallocate(m_device_data); + } + } + + /*! + * \brief Copy assignment operator. + */ + CopyHidingArrayManager& operator=(const CopyHidingArrayManager& other) + { + if (this != &other) + { + // Copy-assign or copy members + m_host_allocator = other.m_host_allocator; + m_device_allocator = other.m_device_allocator; + m_touch = other.m_touch; + + // Allocate new resources before releasing old ones for strong exception safety + void* new_host_data = nullptr; + void* new_device_data = nullptr; + + if (other.m_host_data) + { + new_host_data = m_host_allocator.allocate(other.m_size); + m_resource_manager.copy(new_host_data, other.m_host_data, other.m_size*sizeof(ElementT)); + } + + if (other.m_device_data) + { + new_device_data = m_device_allocator.allocate(other.m_size); + m_resource_manager.copy(new_device_data, other.m_device_data, other.m_size*sizeof(ElementT)); + } + + // Clean up old resources + if (m_host_data) + { + m_host_allocator.deallocate(m_host_data); + } + + if (m_device_data) + { + m_device_allocator.deallocate(m_device_data); + } + + // Assign new resources and size + m_host_data = new_host_data; + m_device_data = new_device_data; + m_size = other.m_size; + } + + return *this; + } + + /*! + * \brief Move assignment operator. + */ + CopyHidingArrayManager& operator=(CopyHidingArrayManager&& other) noexcept + { + if (this != &other) + { + // Release any resources currently held + if (m_host_data) + { + m_host_allocator.deallocate(m_host_data); + m_host_data = nullptr; + } + if (m_device_data) + { + m_device_allocator.deallocate(m_device_data); + m_device_data = nullptr; + } + + // Move-assign or copy members + m_host_allocator = other.m_host_allocator; + m_device_allocator = other.m_device_allocator; + m_size = other.m_size; + m_host_data = other.m_host_data; + m_device_data = other.m_device_data; + m_touch = other.m_touch; + + // Null out other's pointers and reset size + other.m_host_data = nullptr; + other.m_device_data = nullptr; + other.m_size = 0; + other.m_touch = NONE; + } + return *this; + } + + /*! + * \brief Resize the underlying arrays. + */ + virtual void resize(std::size_t newSize) override + { + if (newSize != m_size) + { + if (m_touch == HOST) + { + m_resource_manager.reallocate(m_host_data, newSize); + + if (m_device_data) + { + m_resource_manager.deallocate(m_device_data); + m_device_data = m_device_allocator.allocate(newSize); + } + } + else if (m_touch == DEVICE) + { + m_resource_manager.reallocate(m_device_data, newSize); + + if (m_host_data) + { + m_resource_manager.deallocate(m_host_data); + m_host_data = m_host_allocator.allocate(newSize); + } + } + else + { + if (m_device_data) + { + m_resource_manager.reallocate(m_device_data, newSize); + } + + if (m_host_data) + { + m_resource_manager.reallocate(m_host_data, newSize); + } + } + m_size = newSize; + } + } + + /*! + * \brief Get the size of the underlying arrays. + */ + virtual std::size_t size() const override + { + return m_size; + } + + /*! + * \brief Updates the data to be coherent in the current execution space. + */ + virtual ElementT* data(Context context, bool touch) override + { + if (context == HOST) + { + if (!m_host_data) + { + m_host_data = static_cast(m_host_allocator.allocate(m_size*sizeof(ElementT))); + } + + if (m_touch == DEVICE) + { + m_resource_manager.copy(m_host_data, m_device_data, m_size*sizeof(ElementT)); + m_touch = NONE; + } + + if (touch) + { + m_touch = HOST; + } + + return m_host_data; + } + else if (context == DEVICE) + { + if (!m_device_data) + { + m_device_data = m_device_allocator.allocate(m_size); + } + + if (m_touch == HOST) + { + m_resource_manager.copy(m_device_data, m_host_data, m_size*sizeof(ElementT)); + m_touch = NONE; + } + + if (touch) + { + m_touch = DEVICE; + } + + return m_device_data; + } + else + { + return nullptr; + } + } + + /*! + * \brief Returns the value at index i. + * + * Note: Use this function sparingly as it may be slow. + * + * \param i The index of the element to get. + * \return The value at index i. + */ + virtual ElementT get(std::size_t i) const override { + ElementT result; + + if (m_touch == HOST) { + return m_host_data[i]; + } + else if (m_touch == DEVICE) { + m_resource_manager.copy(&result, m_device_data + i, sizeof(ElementT)); + } + else { + if (m_host_data) { + return m_host_data[i]; + } + else { + throw std::runtime_exception("Reading uninitialized memory"); + } + } + + return ElementT{}; + } + + /*! + * \brief Sets the value at index i to the specified value. + * + * Note: Use this function sparingly as it may be slow. + * + * \param i The index of the element to set. + * \param value The value to set at index i. + */ + virtual void set(std::size_t i, const ElementT& value) override + { + if (m_touch == HOST) { + m_host_data[i] = value; + } + else if (m_touch == DEVICE) { + m_resource_manager.copy(m_device_data + i, &value, sizeof(ElementT)); + } + else { + if (m_host_data) { + m_host_data[i] = value; + } + + if (m_device_data) { + m_resource_manager.copy(m_device_data + i, &value, sizeof(ElementT)); + } + } + } + + private: + umpire::ResourceManager& m_resource_manager{umpire::ResourceManager::getInstance()}; + umpire::Allocator m_host_allocator{m_resource_manager.getAllocator("HOST")}; + umpire::Allocator m_device_allocator{m_resource_manager.getAllocator("DEVICE")}; + std::size_t m_size{0}; + ElementT* m_host_data{nullptr}; + ElementT* m_device_data{nullptr}; + ExecutionContext m_touch{NONE}; + }; // class CopyHidingArrayManager +} // namespace expt +} // namespace chai + +#endif // CHAI_COPY_HIDING_ARRAY_MANAGER_HPP diff --git a/src/chai/expt/DiscreteMemoryArray.hpp b/src/chai/expt/DiscreteMemoryArray.hpp new file mode 100644 index 00000000..4d76a436 --- /dev/null +++ b/src/chai/expt/DiscreteMemoryArray.hpp @@ -0,0 +1,202 @@ +#ifndef CHAI_UNIFIED_MEMORY_ARRAY_HPP +#define CHAI_UNIFIED_MEMORY_ARRAY_HPP + +#include "chai/expt/ExecutionContext.hpp" +#include "umpire/ResourceManager.hpp" + +namespace chai { +namespace expt { + /*! + * \class UnifiedMemoryArray + * + * \brief A container for managing the lifetime and coherence of a + * unified memory array, meaning an array with a single address + * that is accessible from all processors/devices in a system. + * + * This container should be used in tandem with the ExecutionContextManager. + * Together, they provide a programming model where work (e.g. a kernel) + * is generally performed asynchronously, with synchronization occurring + * only as needed for coherence of the array. For example, if the array is + * written to in an asynchronize kernel on a GPU, then the GPU will be + * synchronized if the array needs to be accessed on the CPU. + * + * This model works well for APUs where the CPU and GPU have the same + * physical memory. It also works for pinned (i.e. page-locked) memory + * and in some cases for pageable memory, though no pre-fetching is + * performed. + * + * Example: + * + * \code + * // Create a UnifiedMemoryArray with size 100 and default allocator + * int size = 10000; + * UnifiedMemoryArray array(size); + * + * // Access elements on the device + * std::span device_view(array.data(ExecutionContext::DEVICE, array.size()); + * + * // Launch a kernel that modifies device_view. + * // Note that this example relies on c++20 and the ability to use constexpr + * // host code on the device. + * + * // Access elements on the host. This will synchronize the device. + * std::span host_view(array.data(ExecutionContext::HOST), array.size()); + * + * for (int i = 0; i < size; ++i) { + * std::cout << host_view[i] << "\n"; + * } + * + * // Access and modify individual elements in the container. + * // This should be used sparingly or it will tank performance. + * // Getting the last element after performing a scan is one use case. + * array.get(ExecutionContext::HOST, size - 1) = 10; + * \endcode + */ + template + class UnifiedMemoryArray + public: + UnifiedMemoryArray() = default; + + explicit UnifiedMemoryArray(const umpire::Allocator& allocator) : + m_allocator{allocator} + { + } + + UnifiedMemoryArray(std::size_t size, const umpire::Allocator& allocator) : + m_size{size}, + m_allocator{allocator} + { + m_data = m_allocator.allocate(m_size * sizeof(T)); + // TODO: Investigate if/when to do initialization + } + + explicit UnifiedMemoryArray(int allocatorID) : + m_allocator{umpire::ResourceManager::getInstance().getAllocator(allocatorID)} + { + } + + UnifiedMemoryArray(std::size_t size, int allocatorID) : + m_size{size}, + m_allocator{umpire::ResourceManager::getInstance().getAllocator(allocatorID)} + { + m_data = m_allocator.allocate(m_size * sizeof(T)); + // TODO: Investigate if/when to do initialization + } + + UnifiedMemoryArray(const UnifiedMemoryArray& other) : + m_size{other.m_size}, + m_allocator{other.m_allocator} + { + m_data = m_allocator.allocate(m_size * sizeof(T)); + ExecutionContextManager::getInstance().setExecutionContext(ExecutionContext::DEVICE); + umpire::ResourceManager::getInstance().copy(other.m_data, m_data, m_size * sizeof(T)); + ExecutionContextManager::getInstance().setExecutionContext(ExecutionContext::NONE); + m_last_execution_context = ExecutionContext::DEVICE; + } + + UnifiedMemoryArray(UnifiedMemoryArray&& other) : + m_data{other.m_data}, + m_size{other.m_size}, + m_last_execution_context{other.m_last_execution_context}, + m_allocator{other.m_allocator} + { + other.m_data = nullptr; + other.m_size = 0; + other.m_last_execution_context = NONE; + other.m_allocator = umpire::Allocator(); + } + + UnifiedMemoryArray& operator=(const UnifiedMemoryArray& other) { + if (&other != this) { // Prevent self-assignment + m_allocator.deallocate(m_data); + + m_size = other.m_size; + m_allocator = other.m_allocator; + m_data = m_allocator.allocate(m_size * sizeof(T)); + ExecutionContextManager::getInstance().setExecutionContext(ExecutionContext::DEVICE); + umpire::ResourceManager::getInstance().copy(other.m_data, m_data, m_size * sizeof(T)); + ExecutionContextManager::getInstance().setExecutionContext(ExecutionContext::NONE); + m_last_execution_context = ExecutionContext::DEVICE; + } + + return *this; + } + + UnifiedMemoryArray& operator=(UnifiedMemoryArray&& other) { + if (&other != this) { + m_allocator.deallocate(m_data); + + m_data = other.m_data; + m_size = other.m_size; + m_last_execution_context = other.m_last_execution_context; + m_allocator = other.m_allocator; + + other.m_data = nullptr; + other.m_size = 0; + other.m_last_execution_context = ExecutionContext::NONE; + other.m_allocator = umpire::Allocator(); + } + + return *this; + } + + /*! + * \brief Destructor. + */ + ~UnifiedMemoryArray() { + m_allocator.deallocate(m_data); + } + + /*! + * \brief Get the number of elements. + */ + size_t size() const { + return m_size; + } + + T* data(ExecutionContext executionContext) { + if (executionContext != m_last_execution_context) { + ExecutionContextManager::getInstance().synchronize(m_last_execution_context); + m_last_execution_context = executionContext; + } + + return m_data; + } + + const T* data(ExecutionContext executionContext) const { + if (executionContext != m_last_execution_context) { + ExecutionContextManager::getInstance().synchronize(m_last_execution_context); + m_last_execution_context = ExecutionContext::NONE; + } + + return m_data; + } + + T& get(ExecutionContext executionContext, size_t i) { + if (executionContext != m_last_execution_context) { + ExecutionContextManager::getInstance().synchronize(m_last_execution_context); + m_last_execution_context = executionContext; + } + + return m_data[i]; + } + + const T& get(ExecutionContext executionContext, size_t i) { + if (executionContext != m_last_execution_context) { + ExecutionContextManager::getInstance().synchronize(m_last_execution_context); + m_last_execution_context = ExecutionContext::NONE; + } + + return m_data[i]; + } + + private: + T* m_data{nullptr}; + size_t m_size{0}; + ExecutionContext m_last_execution_context{ExecutionContext::NONE}; + umpire::Allocator m_allocator{}; + }; // class UnifiedMemoryArray +} // namespace expt +} // namespace chai + +#endif // CHAI_UNIFIED_MEMORY_ARRAY_HPP diff --git a/src/chai/expt/ExecutionContext.hpp b/src/chai/expt/ExecutionContext.hpp new file mode 100644 index 00000000..44819004 --- /dev/null +++ b/src/chai/expt/ExecutionContext.hpp @@ -0,0 +1,28 @@ +////////////////////////////////////////////////////////////////////////////// +// Copyright (c) 2016-25, Lawrence Livermore National Security, LLC and CHAI +// project contributors. See the CHAI LICENSE file for details. +// +// SPDX-License-Identifier: BSD-3-Clause +////////////////////////////////////////////////////////////////////////////// + +#ifndef CHAI_EXECUTION_CONTEXT_HPP +#define CHAI_EXECUTION_CONTEXT_HPP + +namespace chai { +namespace expt { + /*! + * \enum ExecutionContext + * + * \brief Represents the state of a program. ArrayManagers update coherence based on the context. + */ + enum ExecutionContext { + NONE = 0, ///< Represents no context. + HOST ///< Represents the host context (i.e. the CPU). +#if defined(CHAI_ENABLE_CUDA) || defined(CHAI_ENABLE_HIP) || defined(CHAI_ENABLE_GPU_SIMULATION_MODE) + , DEVICE ///< Represents the device context (i.e. the GPU). +#endif + }; +} // namespace expt +} // namespace chai + +#endif // CHAI_EXECUTION_CONTEXT_HPP \ No newline at end of file diff --git a/src/chai/expt/ExecutionContextGuard.hpp b/src/chai/expt/ExecutionContextGuard.hpp new file mode 100644 index 00000000..cfaab279 --- /dev/null +++ b/src/chai/expt/ExecutionContextGuard.hpp @@ -0,0 +1,33 @@ +////////////////////////////////////////////////////////////////////////////// +// Copyright (c) 2016-25, Lawrence Livermore National Security, LLC and CHAI +// project contributors. See the CHAI LICENSE file for details. +// +// SPDX-License-Identifier: BSD-3-Clause +////////////////////////////////////////////////////////////////////////////// + +#include "chai/expt/ExecutionContext.hpp" +#include "chai/expt/ExectuionContextManager.hpp" + +#ifndef CHAI_EXECUTION_CONTEXT_GUARD_HPP +#define CHAI_EXECUTION_CONTEXT_GUARD_HPP + +namespace chai { +namespace expt { + class ExecutionContextGuard { + public: + explicit ExecutionContextGuard(ExecutionContext executionContext) { + m_execution_context_manager.setExecutionContext(executionContext); + } + + ~ExecutionContextGuard() { + m_execution_context_manager.setExecutionContext(m_last_execution_context); + } + + private: + ExecutionContextManager& m_execution_context_manager{ExecutionContextManager::getInstance()}; + ExecutionContext m_last_execution_context{m_execution_context_manager.getExecutionContext()}; + }; +} // namespace expt +} // namespace chai + +#endif // CHAI_EXECUTION_CONTEXT_GUARD_HPP \ No newline at end of file diff --git a/src/chai/expt/ExecutionContextManager.hpp b/src/chai/expt/ExecutionContextManager.hpp new file mode 100644 index 00000000..e6f5abec --- /dev/null +++ b/src/chai/expt/ExecutionContextManager.hpp @@ -0,0 +1,145 @@ +////////////////////////////////////////////////////////////////////////////// +// Copyright (c) 2016-25, Lawrence Livermore National Security, LLC and CHAI +// project contributors. See the CHAI LICENSE file for details. +// +// SPDX-License-Identifier: BSD-3-Clause +////////////////////////////////////////////////////////////////////////////// + +#ifndef CHAI_EXECUTION_CONTEXT_MANAGER_HPP +#define CHAI_EXECUTION_CONTEXT_MANAGER_HPP + +#include "chai/expt/ExecutionContext.hpp" + +namespace chai { +namespace expt { + /*! + * \class ExecutionContextManager + * + * \brief Singleton class for managing the current execution context. + * + * This class provides a centralized way to get and set the current execution + * context across the application. + */ + class ExecutionContextManager { + public: + /*! + * \brief Get the singleton instance of ExecutionContextManager. + * + * \return The singleton instance. + */ + static ExecutionContextManager& getInstance() { + static ExecutionContextManager s_instance; + return s_instance; + } + + /*! + * \brief Deleted copy constructor to prevent copying. + */ + ExecutionContextManager(const ExecutionContextManager&) = delete; + + /*! + * \brief Deleted assignment operator to prevent assignment. + */ + ExecutionContextManager& operator=(const ExecutionContextManager&) = delete; + + /*! + * \brief Get the current execution context. + * + * \return The current context. + */ + ExecutionContext getExecutionContext() const { + return m_execution_context; + } + + /*! + * \brief Set the current execution context. + * + * \param context The new context to set. + */ + void setExecutionContext(ExecutionContext context) { + m_execution_context = context; + m_synchronized[context] = false; + } + + /*! + * \brief Synchronize the given execution context. + * + * \param context The execution context that needs synchronization. + */ + void synchronize(ExecutionContext context) { + auto it = m_synchronized.find(context); + + if (it != m_synchronized.end()) { + #if defined(CHAI_ENABLE_DEVICE) + if (context == ExecutionContext::DEVICE) { +#if defined(CHAI_ENABLE_CUDA) + cudaDeviceSynchronize(); +#elif defined(CHAI_ENABLE_HIP) + hipDeviceSynchronize(); +#endif + } + } + bool& unsynchronized = m_unsynchronized[context]; + + if (unsynchronized) { +#if defined(CHAI_ENABLE_DEVICE) + if (context == ExecutionContext::DEVICE) { +#if defined(CHAI_ENABLE_CUDA) + cudaDeviceSynchronize(); +#elif defined(CHAI_ENABLE_HIP) + hipDeviceSynchronize(); +#endif + } + + unsynchronized = false; + } + } + + /*! + * \brief Check if a specific execution context needs synchronization. + * + * \param context The execution context to check. + * \return True if the context needs synchronization, false otherwise. + */ + bool isSynchronized(ExecutionContext context) const { + auto it = m_synchronized.find(context); + + if (it == m_synchronized.end()) { + return true; + } + else { + return it->second; + } + } + + /*! + * \brief Mark the given execution context as synchronized. + * + * This should only be called after synchronization has been performed. + * + * \param context The execution context to clear the synchronization flag for. + */ + void markSynchronized(ExecutionContext context) { + m_synchronized[context] = true; + } + + private: + /*! + * \brief Private constructor for singleton pattern. + */ + constexpr ExecutionContextManager() noexcept = default; + + /*! + * \brief The current execution context. + */ + ExecutionContext m_execution_context = ExecutionContext::NONE; + + /*! + * \brief Map for tracking which execution contexts are synchronized. + */ + std::unordered_map m_synchronized; + }; // class ExecutionContextManager +} // namespace expt +} // namespace chai + +#endif // CHAI_EXECUTION_CONTEXT_MANAGER_HPP diff --git a/src/chai/expt/HostArray.hpp b/src/chai/expt/HostArray.hpp new file mode 100644 index 00000000..8ca03c31 --- /dev/null +++ b/src/chai/expt/HostArray.hpp @@ -0,0 +1,279 @@ +#ifndef CHAI_HOST_ARRAY_HPP +#define CHAI_HOST_ARRAY_HPP + +#include "umpire/ResourceManager.hpp" +#include // for size_t + +namespace chai { +namespace expt { + /*! + * \class HostArray + * + * \brief Manages an array in host memory with RAII semantics. + * + * This class provides a simpler alternative to CopyHidingArray + * when only host memory access is needed. + */ + template + class HostArray { + public: + using size_type = std::size_t; + + /*! + * \brief Default constructor creates an empty array. + */ + HostArray() = default; + + /*! + * \brief Constructs a HostArray with the given Umpire allocator. + */ + HostArray(const umpire::Allocator& allocator) : + m_allocator{allocator} + { + } + + /*! + * \brief Constructs a HostArray with the given Umpire allocator ID. + */ + HostArray(int allocatorID) : + m_resource_manager{umpire::ResourceManager::getInstance()}, + m_allocator{m_resource_manager.getAllocator(allocatorID)} + { + } + + /*! + * \brief Constructs a HostArray with the given size using the default allocator. + */ + HostArray(size_type size) : + m_size{size} + { + if (size > 0) { + m_data = static_cast(m_allocator.allocate(size * sizeof(ElementT))); + } + } + + /*! + * \brief Constructs a HostArray with the given size using the specified allocator. + */ + HostArray(size_type size, const umpire::Allocator& allocator) : + m_allocator{allocator}, + m_size{size} + { + if (size > 0) { + m_data = static_cast(m_allocator.allocate(size * sizeof(ElementT))); + } + } + + /*! + * \brief Constructs a HostArray with the given size using the specified allocator ID. + */ + HostArray(size_type size, int allocatorID) : + m_resource_manager{umpire::ResourceManager::getInstance()}, + m_allocator{m_resource_manager.getAllocator(allocatorID)}, + m_size{size} + { + if (size > 0) { + m_data = static_cast(m_allocator.allocate(size * sizeof(ElementT))); + } + } + + /*! + * \brief Constructs a deep copy of the given HostArray. + */ + HostArray(const HostArray& other) : + m_allocator{other.m_allocator}, + m_size{other.m_size} + { + if (m_size > 0) { + m_data = static_cast(m_allocator.allocate(m_size * sizeof(ElementT))); + + for (size_type i = 0; i < m_size; ++i) { + m_data[i] = other.m_data[i]; + } + } + } + + /*! + * \brief Constructs a HostArray that takes ownership of the resources from the given HostArray. + */ + HostArray(HostArray&& other) noexcept : + m_allocator{other.m_allocator}, + m_size{other.m_size}, + m_data{other.m_data} + { + other.m_size = 0; + other.m_data = nullptr; + } + + /*! + * \brief Destructor releases allocated memory. + */ + ~HostArray() + { + if (m_data) { + m_allocator.deallocate(m_data); + } + } + + /*! + * \brief Copy assignment operator. + */ + HostArray& operator=(const HostArray& other) + { + if (this != &other) { + // Allocate new data before releasing old data + ElementT* new_data = nullptr; + + if (other.m_size > 0) { + new_data = static_cast(other.m_allocator.allocate(other.m_size * sizeof(ElementT))); + + for (size_type i = 0; i < other.m_size; ++i) { + new_data[i] = other.m_data[i]; + } + } + + // Clean up old data + if (m_data) { + m_allocator.deallocate(m_data); + } + + // Update allocator, size, and data + m_allocator = other.m_allocator; + m_size = other.m_size; + m_data = new_data; + } + + return *this; + } + + /*! + * \brief Move assignment operator. + */ + HostArray& operator=(HostArray&& other) noexcept + { + if (this != &other) { + // Clean up current resources + if (m_data) { + m_allocator.deallocate(m_data); + } + + // Move resources from other + m_allocator = other.m_allocator; + m_size = other.m_size; + m_data = other.m_data; + + // Reset other + other.m_size = 0; + other.m_data = nullptr; + } + + return *this; + } + + /*! + * \brief Resize the array. + * + * If the new size is larger, the existing content is preserved and + * new elements are default-initialized. If the new size is smaller, + * only the first newSize elements are preserved. + */ + void resize(size_type newSize) + { + if (newSize != m_size) { + ElementT* new_data = nullptr; + + if (newSize > 0) { + new_data = static_cast(m_allocator.allocate(newSize * sizeof(ElementT))); + + // Copy existing data, up to the smaller of m_size and newSize + size_type copy_size = (newSize < m_size) ? newSize : m_size; + for (size_type i = 0; i < copy_size; ++i) { + new_data[i] = m_data[i]; + } + + // Initialize new elements if expanding + for (size_type i = m_size; i < newSize; ++i) { + new_data[i] = ElementT(); + } + } + + // Clean up old data + if (m_data) { + m_allocator.deallocate(m_data); + } + + m_data = new_data; + m_size = newSize; + } + } + + /*! + * \brief Get the size of the array. + */ + size_type size() const + { + return m_size; + } + + /*! + * \brief Get access to the underlying data. + */ + ElementT* data() + { + return m_data; + } + + /*! + * \brief Get const access to the underlying data. + */ + const ElementT* data() const + { + return m_data; + } + + /*! + * \brief Array subscript operator for element access. + */ + ElementT& operator[](size_type index) + { + return m_data[index]; + } + + /*! + * \brief Const array subscript operator for element access. + */ + const ElementT& operator[](size_type index) const + { + return m_data[index]; + } + + /*! + * \brief Check if the array is empty. + */ + bool empty() const + { + return m_size == 0; + } + + /*! + * \brief Clear the array by deallocating memory and setting size to 0. + */ + void clear() + { + if (m_data) { + m_allocator.deallocate(m_data); + m_data = nullptr; + } + m_size = 0; + } + + private: + umpire::ResourceManager& m_resource_manager{umpire::ResourceManager::getInstance()}; + umpire::Allocator m_allocator{m_resource_manager.getAllocator("HOST")}; + size_type m_size{0}; + ElementT* m_data{nullptr}; + }; // class HostArray +} // namespace expt +} // namespace chai + +#endif // CHAI_HOST_ARRAY_HPP \ No newline at end of file diff --git a/src/chai/expt/HostManager.cpp b/src/chai/expt/HostManager.cpp new file mode 100644 index 00000000..68053494 --- /dev/null +++ b/src/chai/expt/HostManager.cpp @@ -0,0 +1,35 @@ +#include "chai/expt/HostManager.hpp" +#include "umpire/ResourceManager.hpp" + +namespace chai { +namespace expt { + HostManager::HostManager(int allocatorID, std::size_t size) : + Manager{}, + m_allocator_id{allocatorID}, + m_size{size} + { + m_data = umpire::ResourceManager::getInstance().getAllocator(m_allocator_id).allocate(size); + } + + HostManager::~HostManager() { + umpire::ResourceManager::getInstance().getAllocator(m_allocator_id).deallocate(m_data); + } + + std::size_t HostManager::size() const { + return m_size; + } + + void* HostManager::data(ExecutionContext context, bool /* touch */) { + if (context == ExecutionContext::HOST) { + return m_data; + } + else { + return nullptr; + } + } + + int HostManager::getAllocatorID() const { + return m_allocator_id; + } +} // namespace expt +} // namespace chai diff --git a/src/chai/expt/HostManager.hpp b/src/chai/expt/HostManager.hpp new file mode 100644 index 00000000..51fd9eeb --- /dev/null +++ b/src/chai/expt/HostManager.hpp @@ -0,0 +1,60 @@ +#ifndef CHAI_HOST_MANAGER_HPP +#define CHAI_HOST_MANAGER_HPP + +#include "chai/expt/Manager.hpp" + +namespace chai { +namespace expt { + /*! + * \class HostManager + * + * \brief Controls the coherence of an array on the CPU. + */ + class HostManager : public Manager { + public: + /*! + * \brief Constructs a host array manager. + */ + HostManager(int allocatorID, std::size_t size); + + /*! + * \brief Copy constructor is deleted. + */ + HostManager(const HostManager&) = delete; + + /*! + * \brief Copy assignment operator is deleted. + */ + HostManager& operator=(const HostManager&) = delete; + + /*! + * \brief Virtual destructor. + */ + virtual ~HostManager(); + + /*! + * \brief Get the number of elements. + */ + virtual std::size_t size() const override; + + /*! + * \brief Updates the data to be coherent in the current execution space. + * + * \param data [out] A coherent array in the current execution space. + */ + virtual void* data(ExecutionContext context, bool touch) override; + + /*! + * \brief Get the allocator ID. + */ + int getAllocatorID() const; + + private: + int m_allocator_id{-1}; + std::size_t m_size{0}; + void* m_data{nullptr}; + }; // class HostManager +} // namespace expt +} // namespace chai + +#endif // CHAI_HOST_MANAGER_HPP diff --git a/src/chai/expt/ManagedArray.hpp b/src/chai/expt/ManagedArray.hpp new file mode 100644 index 00000000..2df313cb --- /dev/null +++ b/src/chai/expt/ManagedArray.hpp @@ -0,0 +1,251 @@ +#ifndef CHAI_MANAGED_ARRAY_HPP +#define CHAI_MANAGED_ARRAY_HPP + +#include "chai/expt/ArrayManager.hpp" +#include + +namespace chai { +namespace expt { + /*! + * \class ManagedArray + * + * \brief An array class that manages coherency across the CPU and GPU. + * How the coherence is obtained is controlled by the array manager. + * + * \tparam ElementType The type of element in the array. + */ + template + class ManagedArray { + public: + /*! + * \brief Constructs an empty array without an array manager. + */ + ManagedArray() = default; + + /*! + * \brief Constructs an array from a manager. + * + * \param manager The array manager controls the coherence of the array. + * + * \note The array takes ownership of the manager. + */ + explicit ManagedArray(ArrayManager* manager) + : m_manager{manager} + {} + + /*! + * \brief Constructs a shallow copy of an array from another and makes + * the data coherent in the current execution space. + * + * \param other The other array. + * + * \note This is a shallow copy. + */ + CHAI_HOST_DEVICE ManagedArray(const ManagedArray& other) : + m_data{other.m_data}, + m_size{other.m_size}, + m_manager{other.m_manager} + { + update(); + } + + /*! + * \brief Constructs a ManagedArray from a ManagedArray. + * + * \param other The non-const array to convert from. + * + * \note This is a converting constructor that enables implicit conversion + * from ManagedArray to ManagedArray. + */ + template * = nullptr> + CHAI_HOST_DEVICE ManagedArray(const ManagedArray& other) + : m_data{other.m_data}, + m_size{other.m_size}, + m_manager{other.m_manager} + { + } + + /*! + * \brief Sets the array manager for this ManagedArray. + * + * \param manager The new array manager to be set. + * + * \post The ManagedArray takes ownership of the new manager objet. + */ + void setManager(ArrayManager* manager) + { + delete m_manager; + m_manager = manager; + } + + /*! + * \brief Get the array manager associated with this ManagedArray. + * + * \return A pointer to the array manager. + */ + ArrayManager* getManager() const { + return m_manager; + } + + /*! + * \brief Resizes the array to the specified new size. + * + * \param newSize The new size to resize the array to. + * + * \note This method updates the size of the array and triggers a resize operation in the array manager if it exists. + * If no array manager is associated, an exception is thrown. + */ + void resize(std::size_t newSize) { + if (m_manager) { + m_size = newSize; + m_manager->resize(newSize); + } + else { + throw std::runtime_exception("Unable to resize"); + } + } + + /*! + * \brief Frees the resources associated with this array. + * + * \note Once free has been called, it is invalid to use any other copies + * of this array (since copies are shallow). + */ + void free() { + m_data = nullptr; + m_size = 0; + delete m_manager; + m_manager = nullptr; + } + + /*! + * \brief Get the number of elements in the array. + * + * \pre The copy constructor has been called with the execution space + * set to CPU or GPU (e.g. by the RAJA plugin). + */ + CHAI_HOST_DEVICE std::size_t size() const { + return m_size; + } + + CHAI_HOST_DEVICE void update() const { +#if !defined(CHAI_DEVICE_COMPILE) + if (m_manager) { + m_data = m_manager->data(!std::is_const_v); + } +#endif + } + + CHAI_HOST_DEVICE void cupdate() const { +#if !defined(CHAI_DEVICE_COMPILE) + if (m_manager) { + m_data = m_manager->data(false); + } +#endif + } + + /*! + * \brief Get a pointer to the element data in the specified context. + * + * \param context The context in which to retrieve the element data. + * + * \return A pointer to the element data in the specified context. + */ + CHAI_HOST_DEVICE ElementType* data() const { + update(); + return m_data; + } + + /*! + * \brief Get a const pointer to the element data in the specified context. + * + * \param context The context in which to retrieve the const element data. + * + * \return A const pointer to the element data in the specified context. + */ + CHAI_HOST_DEVICE const ElementType* cdata() const { + cupdate(); + return m_data; + } + + /*! + * \brief Get the ith element in the array. + * + * \param i The index of the element to retrieve. + * + * \pre The copy constructor has been called with the execution space + * set to CPU or GPU (e.g. by the RAJA plugin). + */ + CHAI_HOST_DEVICE ElementType& operator[](std::size_t i) const { + return m_data[i]; + } + + /*! + * \brief Get the value of the element at the specified index. + * + * \param i The index of the element to retrieve. + * + * \return The value of the element at the specified index. + * + * \throw std::runtime_exception if unable to retrieve the element. + */ + ElementType get(std::size_t i) const { + if (m_manager) { + return *(static_cast(m_manager->get(i*sizeof(ElementType), sizeof(ElementType)))); + } + else { + throw std::runtime_exception("Unable to get element"); + } + } + + /*! + * \brief Set a value at a specified index in the array. + * + * \param i The index where the value is to be set. + * \param value The value to set at the specified index. + * + * \throw std::runtime_exception if the array manager is not associated with the ManagedArray. + */ + void set(std::size_t i, const ElementType& value) { + if (m_manager) { + m_manager->set(i*sizeof(ElementType), sizeof(ElementType), static_cast(std::addressof(value))); + } + else { + throw std::runtime_exception("Unable to set element"); + } + } + + private: + /*! + * The array that is coherent in the current execution space. + */ + ElementType* m_data = nullptr; + + /*! + * The number of elements in the array. + */ + std::size_t m_size = 0; + + /*! + * The array manager controls the coherence of the array. + */ + ArrayManager* m_manager = nullptr; + }; // class ManagedArray + + /*! + * \brief Constructs an array by creating a new manager object. + * + * \tparam ArrayManager The type of array manager. + * \tparam Args The type of the arguments used to construct the array manager. + * + * \param args The arguments to construct an array manager. + */ + template + ManagedArray makeArray(Args&&... args) { + return ManagedArray(new ArrayManager(std::forward(args)...)); + } +} // namespace expt +} // namespace chai + +#endif // CHAI_MANAGED_ARRAY_HPP diff --git a/src/chai/expt/ManagedArrayView.hpp b/src/chai/expt/ManagedArrayView.hpp new file mode 100644 index 00000000..c6d31ffa --- /dev/null +++ b/src/chai/expt/ManagedArrayView.hpp @@ -0,0 +1,151 @@ +#ifndef CHAI_ARRAY_VIEW_HPP +#define CHAI_ARRAY_VIEW_HPP + +#include "chai/Manager.hpp" +#include + +namespace chai { +namespace expt { + /*! + * \class ArrayView + * + * \brief A view into an existing Array without taking ownership of the data. + * + * \tparam T The type of element in the array view. + */ + template + class ArrayView { + public: + /*! + * \brief Constructs an empty array view. + */ + ArrayView() = default; + + /*! + * \brief Constructs an array view from a manager. + * + * \param manager The array manager controls the coherence of the array. + */ + explicit ArrayView(Manager* manager) : + m_manager{manager} + { + if (m_manager) + { + m_size = m_manager->size(); + } + } + + /*! + * \brief Constructs an array view with specified size and manager. + * + * \param size The number of elements + * \param manager The array manager + */ + ArrayView(std::size_t offset, std::size_t size, Manager* manager) : + m_offset{offset}, + m_size{size}, + m_manager{manager} + { + } + + /*! + * \brief Constructs a shallow copy of an array view from another and makes + * the data coherent in the current execution space. + * + * \param other The other array view. + * + * \note This is a shallow copy. + */ + CHAI_HOST_DEVICE ArrayView(const ArrayView& other) : + m_data{other.m_data}, + m_offset{other.m_offset}, + m_size{other.m_size}, + m_manager{other.m_manager} + { +#if !defined(CHAI_DEVICE_COMPILE) + if (m_manager) { + m_data = static_cast(m_manager->data(ContextManager::getInstance()::getContext(), !std::is_const::value)) + m_offset; + } +#endif + } + + CHAI_HOST_DEVICE std::size_t offset() const { + return m_offset; + } + + /*! + * \brief Get the number of elements in the array view. + * + * \pre The copy constructor has been called with the execution space + * set to CPU or GPU (e.g. by the RAJA plugin). + */ + CHAI_HOST_DEVICE std::size_t size() const { + return m_size; + } + + CHAI_HOST_DEVICE T* data() const { +#if !defined(CHAI_DEVICE_COMPILE) + if (m_manager) { + m_data = static_cast(m_manager->data(ExecutionContext::HOST, !std::is_const::value)) + m_offset; + } +#endif + return m_data; + } + + CHAI_HOST_DEVICE T* data(ExecutionContext context) const { +#if !defined(CHAI_DEVICE_COMPILE) + if (m_manager) { + m_data = static_cast(m_manager->data(context, !std::is_const::value)) + m_offset; + } +#endif + return m_data; + } + + /*! + * \brief Get the ith element in the array view. + * + * \param i The index of the element to retrieve. + * + * \pre The copy constructor has been called with the execution space + * set to CPU or GPU (e.g. by the RAJA plugin). + */ + CHAI_HOST_DEVICE T& operator[](std::size_t i) const { + return m_data[i]; + } + + private: + /*! + * The array that is coherent in the current execution space. + */ + T* m_data = nullptr; + + /*! + * The starting element in the array view. + */ + std::size_t m_offset = 0; + + /*! + * The number of elements in the array view. + */ + std::size_t m_size = 0; + + /*! + * The array manager controls the coherence of the array. + * ArrayView doesn't own the manager + */ + Manager* m_manager = nullptr; + }; // class ArrayView + + /*! + * \brief Constructs an array view by using an existing manager object. + * + * \tparam Manager The type of array manager. + */ + template + ArrayView makeArrayView(Manager* manager) { + return ArrayView(manager); + } +} // namespace expt +} // namespace chai + +#endif // CHAI_ARRAY_VIEW_HPP \ No newline at end of file diff --git a/src/chai/expt/MemoryManager.hpp b/src/chai/expt/MemoryManager.hpp new file mode 100644 index 00000000..b82e4f24 --- /dev/null +++ b/src/chai/expt/MemoryManager.hpp @@ -0,0 +1,79 @@ +////////////////////////////////////////////////////////////////////////////// +// Copyright (c) 2016-25, Lawrence Livermore National Security, LLC and CHAI +// project contributors. See the CHAI LICENSE file for details. +// +// SPDX-License-Identifier: BSD-3-Clause +////////////////////////////////////////////////////////////////////////////// + +#ifndef CHAI_ARRAY_MANAGER_HPP +#define CHAI_ARRAY_MANAGER_HPP + +#include "chai/expt/Context.hpp" +#include + +namespace chai { +namespace expt { + /*! + * \class ArrayManager + * + * \brief Controls the coherence of an array. + */ + class ArrayManager { + public: + /*! + * \brief Virtual destructor. + */ + virtual ~ArrayManager() = default; + + /*! + * \brief Creates a clone of this ArrayManager. + * + * \return A new ArrayManager object that is a clone of this instance. + */ + virtual ArrayManager* clone() const = 0; + + /*! + * \brief Resizes the array to the specified new size. + * + * \param newSize The new size to resize the array to. + */ + virtual void resize(std::size_t newSize) = 0; + + /*! + * \brief Returns the size of the contained array. + * + * \return The size of the contained array. + */ + virtual std::size_t size() const = 0; + + /*! + * \brief Updates the data to be coherent in the current execution context. + * + * \param data [out] A coherent array in the current execution context. + */ + virtual T* data(Context context, bool touch) = 0; + + /*! + * \brief Returns the value at index i. + * + * Note: Use this function sparingly as it may be slow. + * + * \param i The index of the element to get. + * \return The value at index i. + */ + virtual T get(std::size_t i) const = 0; + + /*! + * \brief Sets the value at index i to the specified value. + * + * Note: Use this function sparingly as it may be slow. + * + * \param i The index of the element to set. + * \param value The value to set at index i. + */ + virtual void set(std::size_t i, const T& value) = 0; + }; // class ArrayManager +} // namespace expt +} // namespace chai + +#endif // CHAI_ARRAY_MANAGER_HPP diff --git a/src/chai/expt/NullArrayManager.hpp b/src/chai/expt/NullArrayManager.hpp new file mode 100644 index 00000000..62d2a42c --- /dev/null +++ b/src/chai/expt/NullArrayManager.hpp @@ -0,0 +1,124 @@ +////////////////////////////////////////////////////////////////////////////// +// Copyright (c) 2016-25, Lawrence Livermore National Security, LLC and CHAI +// project contributors. See the CHAI LICENSE file for details. +// +// SPDX-License-Identifier: BSD-3-Clause +////////////////////////////////////////////////////////////////////////////// + +#ifndef CHAI_NULL_ARRAY_MANAGER_HPP +#define CHAI_NULL_ARRAY_MANAGER_HPP + +#include "chai/expt/ArrayManager.hpp" +#include + +namespace chai { +namespace expt { + +/*! + * \class NullArrayManager + * + * \brief A null implementation of ArrayManager that doesn't actually store data. + * + * This class implements the ArrayManager interface but doesn't actually + * store any data. It's implemented as a singleton. + */ +template +class NullArrayManager : public ArrayManager { +public: + /*! + * \brief Get the singleton instance of NullArrayManager. + * + * \return Reference to the singleton instance. + */ + static NullArrayManager& getInstance() { + static NullArrayManager instance; + return instance; + } + + /*! + * \brief Virtual destructor. + */ + virtual ~NullArrayManager() = default; + + /*! + * \brief Creates a clone of this NullArrayManager. + * + * \return Pointer to the singleton instance. + */ + virtual ArrayManager* clone() const override { + return &getInstance(); + } + + /*! + * \brief Throws an exception when attempting to resize. + * + * \param newSize The new size (ignored). + * \throws std::runtime_error Always throws this exception. + */ + virtual void resize(std::size_t newSize) override { + throw std::runtime_error("Cannot resize NullArrayManager"); + } + + /*! + * \brief Returns 0 as the size. + * + * \return Always returns 0. + */ + virtual std::size_t size() const override { + return 0; + } + + /*! + * \brief Returns a nullptr for data access. + * + * \param context The execution context (ignored). + * \param touch Whether to mark data as touched (ignored). + * \return Always returns nullptr. + */ + virtual T* data(Context context, bool touch) override { + return nullptr; + } + + /*! + * \brief Throws an exception when attempting to get a value. + * + * \param i The index (ignored). + * \return Never returns. + * \throws std::runtime_error Always throws this exception. + */ + virtual T get(std::size_t i) const override { + throw std::runtime_error("Cannot get value from NullArrayManager"); + } + + /*! + * \brief Throws an exception when attempting to set a value. + * + * \param i The index (ignored). + * \param value The value (ignored). + * \throws std::runtime_error Always throws this exception. + */ + virtual void set(std::size_t i, const T& value) override { + throw std::runtime_error("Cannot set value in NullArrayManager"); + } + +private: + /*! + * \brief Private constructor for singleton pattern. + */ + NullArrayManager() {} + + /*! + * \brief Delete copy constructor. + */ + NullArrayManager(const NullArrayManager&) = delete; + + /*! + * \brief Delete assignment operator. + */ + NullArrayManager& operator=(const NullArrayManager&) = delete; +}; + +} // namespace expt +} // namespace chai + +#endif // CHAI_NULL_ARRAY_MANAGER_HPP \ No newline at end of file diff --git a/src/chai/expt/PageableManager.hpp b/src/chai/expt/PageableManager.hpp new file mode 100644 index 00000000..1408b83c --- /dev/null +++ b/src/chai/expt/PageableManager.hpp @@ -0,0 +1,18 @@ +#ifndef CHAI_PAGEABLE_MANAGER_HPP +#define CHAI_PAGEABLE_MANAGER_HPP + +#include "chai/expt/PinnedManager.hpp" + +namespace chai { +namespace expt { + /*! + * \alias PageableManager + * + * \brief Controls the coherence of an array on the host and device. + */ + template + using PageableManager = PinnedManager; +} // namespace expt +} // namespace chai + +#endif // CHAI_PAGEABLE_MANAGER_HPP diff --git a/src/chai/expt/PinnedArrayManager.hpp b/src/chai/expt/PinnedArrayManager.hpp new file mode 100644 index 00000000..a568c5b6 --- /dev/null +++ b/src/chai/expt/PinnedArrayManager.hpp @@ -0,0 +1,295 @@ +////////////////////////////////////////////////////////////////////////////// +// Copyright (c) 2016-25, Lawrence Livermore National Security, LLC and CHAI +// project contributors. See the CHAI LICENSE file for details. +// +// SPDX-License-Identifier: BSD-3-Clause +////////////////////////////////////////////////////////////////////////////// + +#ifndef CHAI_PINNED_ARRAY_MANAGER_HPP +#define CHAI_PINNED_ARRAY_MANAGER_HPP + +#include "chai/expt/ArrayManager.hpp" +#include "chai/expt/ContextManager.hpp" +#include "umpire/ResourceManager.hpp" + +namespace chai { +namespace expt { + /*! + * \class PinnedArrayManager + * + * \brief Controls the coherence of an array on the host and device. + */ + template + class PinnedArrayManager : public ArrayManager { + public: + /*! + * Constructs a PinnedArrayManager with default allocators from Umpire + * for the "HOST" and "DEVICE" resources. + */ + PinnedArrayManager() = default; + + /*! + * Constructs a PinnedArrayManager with the given Umpire allocators. + */ + PinnedArrayManager(const umpire::Allocator& allocator) + : ArrayManager{}, + m_allocator{allocator} + { + } + + /*! + * Constructs a PinnedArrayManager with the given Umpire allocator IDs. + */ + PinnedArrayManager(int allocatorID) + : ArrayManager{}, + m_resource_manager{umpire::ResourceManager::getInstance()}, + m_allocator{m_resource_manager.getAllocator(allocatorID)} + { + } + + /*! + * Constructs a PinnedArrayManager with the given size using default allocators + * from Umpire for the "HOST" and "DEVICE" resources. + */ + PinnedArrayManager(std::size_t size) + : ArrayManager{}, + m_size{size} + { + // TODO: Exception handling + m_data = static_cast(m_allocator.allocate(size*sizeof(ElementT)); + } + + /*! + * Constructs a PinnedArrayManager with the given size using the given Umpire + * allocators. + */ + PinnedArrayManager(std::size_t size, + const umpire::Allocator& allocator) + : ArrayManager{}, + m_allocator{allocator}, + m_size{size} + { + // TODO: Exception handling + m_data = static_cast(m_allocator.allocate(size*sizeof(ElementT)); + } + + /*! + * Constructs a PinnedArrayManager with the given size using the given Umpire + * allocator IDs. + */ + PinnedArrayManager(std::size_t size, + int allocatorID) + : ArrayManager{}, + m_resource_manager{umpire::ResourceManager::getInstance()}, + m_allocator{m_resource_manager.getAllocator(allocatorID)}, + m_size{size} + { + // TODO: Exception handling + static_cast(m_allocator.allocate(size*sizeof(ElementT)); + } + + /*! + * Constructs a deep copy of the given PinnedArrayManager. + */ + PinnedArrayManager(const PinnedArrayManager& other) + : ArrayManager{}, + m_allocator{other.m_allocator}, + m_size{other.m_size}, + m_touch{other.m_touch} + { + if (other.m_data) + { + m_data = m_allocator.allocate(m_size); + m_resource_manager.copy(m_data, other.m_data, m_size*sizeof(ElementT)); + // TODO: The copy could potentially change in which space the last touch occurs + } + } + + /*! + * Constructs a PinnedArrayManager that takes ownership of the + * resources from the given PinnedArrayManager. + */ + PinnedArrayManager(PinnedArrayManager&& other) noexcept + : ArrayManager{}, + m_allocator{other.m_allocator}, + m_size{other.m_size}, + m_touch{other.m_touch}, + m_data{other.m_data} + { + other.m_size = 0; + other.m_data = nullptr; + other.m_touch = NONE; + } + + /*! + * \brief Virtual destructor. + */ + virtual ~PinnedArrayManager() + { + m_allocator.deallocate(m_data); + } + + /*! + * \brief Copy assignment operator. + */ + PinnedArrayManager& operator=(const PinnedArrayManager& other) + { + if (this != &other) + { + // Copy-assign or copy members + m_allocator = other.m_allocator; + m_touch = other.m_touch; + + // Allocate new resources before releasing old ones for strong exception safety + void* new_data = nullptr; + + if (other.m_data) + { + new_data = static_cast(m_allocator.allocate(other.m_size*sizeof(ElementT))); + m_resource_manager.copy(new_data, other.m_data, other.m_size*sizeof(ElementT)); + // TODO: The copy operation could change m_touch + } + + // Clean up old resources + if (m_data) + { + m_allocator.deallocate(m_data); + } + + // Assign new resources and size + m_data = new_data; + m_size = other.m_size; + } + + return *this; + } + + /*! + * \brief Move assignment operator. + */ + PinnedArrayManager& operator=(PinnedArrayManager&& other) noexcept + { + if (this != &other) + { + // Release any resources currently held + if (m_data) + { + m_allocator.deallocate(m_data); + } + + // Move-assign or copy members + m_allocator = other.m_allocator; + m_size = other.m_size; + m_data = other.m_data; + m_touch = other.m_touch; + + // Null out other's pointers and reset size + other.m_data = nullptr; + other.m_size = 0; + other.m_touch = NONE; + } + return *this; + } + + /*! + * \brief Resize the underlying arrays. + */ + virtual void resize(std::size_t newSize) override + { + if (newSize != m_size) + { + // TODO: Is any synchronization needed? + m_resource_manager.reallocate(m_data, newSize); + m_size = newSize; + } + } + + /*! + * \brief Get the size of the underlying arrays. + */ + virtual std::size_t size() const override + { + return m_size; + } + + /*! + * \brief Updates the data to be coherent in the current execution space. + */ + virtual ElementT* data(Context context, bool touch) override + { + ElementT* result{nullptr}; + + if (context == HOST) + { + if (m_touch == DEVICE) + { + m_context_manager.synchronize(DEVICE); + m_touch = NONE; + } + + if (touch) + { + m_touch = HOST; + } + + result = m_data; + } + else if (context == DEVICE) + { + if (m_touch == HOST) + { + // TODO: Should we call m_context_manager.synchronize(HOST)? Would support host openmp. + m_touch = NONE; + } + + if (touch) + { + m_touch = DEVICE; + } + + result = m_data; + } + + return result; + } + + /*! + * \brief Returns the value at index i. + * + * Note: Use this function sparingly as it may be slow. + * + * \param i The index of the element to get. + * \return The value at index i. + */ + virtual ElementT get(std::size_t i) const override { + m_context_manager.synchronize(m_touch); + m_touch = NONE; + return m_data[i]; + } + + /*! + * \brief Sets the value at index i to the specified value. + * + * Note: Use this function sparingly as it may be slow. + * + * \param i The index of the element to set. + * \param value The value to set at index i. + */ + virtual void set(std::size_t i, const ElementT& value) override + { + m_context_manager.synchronize(m_touch); + m_touch = HOST; + m_data[i] = value; + } + + private: + umpire::ResourceManager& m_resource_manager{umpire::ResourceManager::getInstance()}; + umpire::Allocator m_allocator{m_resource_manager.getAllocator("DEVICE")}; + std::size_t m_size{0}; + ElementT* m_data{nullptr}; + ExecutionContext m_touch{NONE}; + }; // class PinnedArrayManager +} // namespace expt +} // namespace chai + +#endif // CHAI_PINNED_ARRAY_MANAGER_HPP diff --git a/src/chai/expt/PinnedArrayView.hpp b/src/chai/expt/PinnedArrayView.hpp new file mode 100644 index 00000000..8a942787 --- /dev/null +++ b/src/chai/expt/PinnedArrayView.hpp @@ -0,0 +1,73 @@ +#ifndef CHAI_PINNED_ARRAY_VIEW_HPP +#define CHAI_PINNED_ARRAY_VIEW_HPP + +#include "chai/expt/PinnedArrayContainer.hpp" + +namespace chai { +namespace expt { + /*! + * \class PinnedArrayView + * + * \brief Provides a non-owning view into a PinnedArrayContainer. + */ + template + class PinnedArrayView { + public: + /*! + * \brief Default constructor. + */ + PinnedArrayView() = default; + + /*! + * \brief Construct a view from a PinnedArrayContainer. + * + * \param container The container to view. + */ + explicit PinnedArrayView(PinnedArrayContainer& container) : + m_container(&container) + { + } + + /*! + * \brief Get the number of elements. + */ + size_t size() const { + return m_container ? m_container->size() : 0; + } + + /*! + * \brief Get pointer to the data for the given execution context. + */ + T* data(ExecutionContext executionContext) { + return m_container ? m_container->data(executionContext) : nullptr; + } + + /*! + * \brief Get const pointer to the data for the given execution context. + */ + const T* data(ExecutionContext executionContext) const { + return m_container ? m_container->data(executionContext) : nullptr; + } + + /*! + * \brief Get element at index i for the given execution context. + */ + T& get(ExecutionContext executionContext, size_t i) { + return m_container->get(executionContext, i); + } + + /*! + * \brief Get const element at index i for the given execution context. + */ + const T& get(ExecutionContext executionContext, size_t i) const { + return m_container->get(executionContext, i); + } + + private: + PinnedArrayContainer* m_container{nullptr}; + }; // class PinnedArrayView + +} // namespace expt +} // namespace chai + +#endif // CHAI_PINNED_ARRAY_VIEW_HPP \ No newline at end of file diff --git a/src/chai/expt/PinnedManager.hpp b/src/chai/expt/PinnedManager.hpp new file mode 100644 index 00000000..2c04a7b1 --- /dev/null +++ b/src/chai/expt/PinnedManager.hpp @@ -0,0 +1,103 @@ +#ifndef CHAI_PINNED_MANAGER_HPP +#define CHAI_PINNED_MANAGER_HPP + +namespace chai { +namespace expt { + /*! + * \class PinnedManager + * + * \brief Controls the coherence of an array on the host and device. + */ + template + class PinnedManager : public Manager { + public: + PinnedManager() noexcept(noexcept(Allocator())) : + PinnedManager(Allocator()) + { + } + + explicit PinnedManager(const Allocator& allocator) : + m_allocator{allocator} + { + } + + /*! + * \brief Constructs a host array manager. + */ + explicit PinnedManager(size_t size, + const Allocator& allocator = Allocator()) : + PinnedManager(allocator), + m_size{size}, + // TODO: Investigate if allocate should be in body of constructor + m_data{allocator.allocate(size)}, + { + } + + PinnedManager(const PinnedManager& other) : + m_size{other.m_size}, + m_data{other.m_allocator.allocate(size)}, + m_allocator{other.m_allocator} + { + // Copy data from other array + } + + PinnedManager(PinnedManager&& other) : + m_size{other.m_size}, + m_data{other.m_data}, + m_allocator{other.m_allocator} + { + other.m_size = 0; + other.m_data = nullptr; + other.m_allocator = Allocator(); + } + + PinnedManager& operator=(const PinnedManager&) = delete; + + /*! + * \brief Virtual destructor. + */ + virtual ~PinnedManager() { + m_allocator.deallocate(m_data); + } + + /*! + * \brief Get the number of elements. + */ + virtual size_t size() const { + return m_size; + } + + virtual T* data(ExecutionContext context, bool touch) { + + } + + /*! + * \brief Updates the data to be coherent in the current execution space. + * + * \param data [out] A coherent array in the current execution space. + */ + virtual void update(void*& data, bool touch) { + ExecutionContext context = execution_context(); + + if (context == ExecutionContext::None) { + data = nullptr; + } + else { + if (context == ExecutionContext::Host) { + // TODO: Only sync if last touched on device + synchronizeDeviceIfNeeded(); + } + + data = m_data; + } + } + + private: + size_t m_size{0}; + T* m_data{nullptr}; + Allocator m_allocator{}; + }; // class PinnedManager +} // namespace expt +} // namespace chai + +#endif // CHAI_PINNED_MANAGER_HPP diff --git a/src/chai/expt/PinnedMemoryManager.hpp b/src/chai/expt/PinnedMemoryManager.hpp new file mode 100644 index 00000000..d14aa85f --- /dev/null +++ b/src/chai/expt/PinnedMemoryManager.hpp @@ -0,0 +1,295 @@ +////////////////////////////////////////////////////////////////////////////// +// Copyright (c) 2016-25, Lawrence Livermore National Security, LLC and CHAI +// project contributors. See the CHAI LICENSE file for details. +// +// SPDX-License-Identifier: BSD-3-Clause +////////////////////////////////////////////////////////////////////////////// + +#ifndef CHAI_PINNED_MEMORY_MANAGER_HPP +#define CHAI_PINNED_MEMORY_MANAGER_HPP + +#include "chai/expt/ArrayManager.hpp" +#include "chai/expt/ContextManager.hpp" +#include "umpire/ResourceManager.hpp" + +namespace chai { +namespace expt { + /*! + * \class PinnedMemoryManager + * + * \brief Controls the coherence of an array on the host and device. + */ + template + class PinnedMemoryManager : public ArrayManager { + public: + /*! + * Constructs a PinnedMemoryManager with default allocators from Umpire + * for the "HOST" and "DEVICE" resources. + */ + PinnedMemoryManager() = default; + + /*! + * Constructs a PinnedMemoryManager with the given Umpire allocators. + */ + PinnedMemoryManager(const umpire::Allocator& allocator) + : ArrayManager{}, + m_allocator{allocator} + { + } + + /*! + * Constructs a PinnedMemoryManager with the given Umpire allocator IDs. + */ + PinnedMemoryManager(int allocatorID) + : ArrayManager{}, + m_resource_manager{umpire::ResourceManager::getInstance()}, + m_allocator{m_resource_manager.getAllocator(allocatorID)} + { + } + + /*! + * Constructs a PinnedMemoryManager with the given size using default allocators + * from Umpire for the "HOST" and "DEVICE" resources. + */ + PinnedMemoryManager(std::size_t size) + : ArrayManager{}, + m_size{size} + { + // TODO: Exception handling + m_data = static_cast(m_allocator.allocate(size*sizeof(ElementT)); + } + + /*! + * Constructs a PinnedMemoryManager with the given size using the given Umpire + * allocators. + */ + PinnedMemoryManager(std::size_t size, + const umpire::Allocator& allocator) + : ArrayManager{}, + m_allocator{allocator}, + m_size{size} + { + // TODO: Exception handling + m_data = static_cast(m_allocator.allocate(size*sizeof(ElementT)); + } + + /*! + * Constructs a PinnedMemoryManager with the given size using the given Umpire + * allocator IDs. + */ + PinnedMemoryManager(std::size_t size, + int allocatorID) + : ArrayManager{}, + m_resource_manager{umpire::ResourceManager::getInstance()}, + m_allocator{m_resource_manager.getAllocator(allocatorID)}, + m_size{size} + { + // TODO: Exception handling + static_cast(m_allocator.allocate(size*sizeof(ElementT)); + } + + /*! + * Constructs a deep copy of the given PinnedMemoryManager. + */ + PinnedMemoryManager(const PinnedMemoryManager& other) + : ArrayManager{}, + m_allocator{other.m_allocator}, + m_size{other.m_size}, + m_touch{other.m_touch} + { + if (other.m_data) + { + m_data = m_allocator.allocate(m_size); + m_resource_manager.copy(m_data, other.m_data, m_size*sizeof(ElementT)); + // TODO: The copy could potentially change in which space the last touch occurs + } + } + + /*! + * Constructs a PinnedMemoryManager that takes ownership of the + * resources from the given PinnedMemoryManager. + */ + PinnedMemoryManager(PinnedMemoryManager&& other) noexcept + : ArrayManager{}, + m_allocator{other.m_allocator}, + m_size{other.m_size}, + m_touch{other.m_touch}, + m_data{other.m_data} + { + other.m_size = 0; + other.m_data = nullptr; + other.m_touch = NONE; + } + + /*! + * \brief Virtual destructor. + */ + virtual ~PinnedMemoryManager() + { + m_allocator.deallocate(m_data); + } + + /*! + * \brief Copy assignment operator. + */ + PinnedMemoryManager& operator=(const PinnedMemoryManager& other) + { + if (this != &other) + { + // Copy-assign or copy members + m_allocator = other.m_allocator; + m_touch = other.m_touch; + + // Allocate new resources before releasing old ones for strong exception safety + void* new_data = nullptr; + + if (other.m_data) + { + new_data = static_cast(m_allocator.allocate(other.m_size*sizeof(ElementT))); + m_resource_manager.copy(new_data, other.m_data, other.m_size*sizeof(ElementT)); + // TODO: The copy operation could change m_touch + } + + // Clean up old resources + if (m_data) + { + m_allocator.deallocate(m_data); + } + + // Assign new resources and size + m_data = new_data; + m_size = other.m_size; + } + + return *this; + } + + /*! + * \brief Move assignment operator. + */ + PinnedMemoryManager& operator=(PinnedMemoryManager&& other) noexcept + { + if (this != &other) + { + // Release any resources currently held + if (m_data) + { + m_allocator.deallocate(m_data); + } + + // Move-assign or copy members + m_allocator = other.m_allocator; + m_size = other.m_size; + m_data = other.m_data; + m_touch = other.m_touch; + + // Null out other's pointers and reset size + other.m_data = nullptr; + other.m_size = 0; + other.m_touch = NONE; + } + return *this; + } + + /*! + * \brief Resize the underlying arrays. + */ + virtual void resize(std::size_t newSize) override + { + if (newSize != m_size) + { + // TODO: Is any synchronization needed? + m_resource_manager.reallocate(m_data, newSize); + m_size = newSize; + } + } + + /*! + * \brief Get the size of the underlying arrays. + */ + virtual std::size_t size() const override + { + return m_size; + } + + /*! + * \brief Updates the data to be coherent in the current execution space. + */ + virtual ElementT* data(Context context, bool touch) override + { + ElementT* result{nullptr}; + + if (context == HOST) + { + if (m_touch == DEVICE) + { + m_context_manager.synchronize(DEVICE); + m_touch = NONE; + } + + if (touch) + { + m_touch = HOST; + } + + result = m_data; + } + else if (context == DEVICE) + { + if (m_touch == HOST) + { + // TODO: Should we call m_context_manager.synchronize(HOST)? Would support host openmp. + m_touch = NONE; + } + + if (touch) + { + m_touch = DEVICE; + } + + result = m_data; + } + + return result; + } + + /*! + * \brief Returns the value at index i. + * + * Note: Use this function sparingly as it may be slow. + * + * \param i The index of the element to get. + * \return The value at index i. + */ + virtual ElementT get(std::size_t i) const override { + m_context_manager.synchronize(m_touch); + m_touch = NONE; + return m_data[i]; + } + + /*! + * \brief Sets the value at index i to the specified value. + * + * Note: Use this function sparingly as it may be slow. + * + * \param i The index of the element to set. + * \param value The value to set at index i. + */ + virtual void set(std::size_t i, const ElementT& value) override + { + m_context_manager.synchronize(m_touch); + m_touch = HOST; + m_data[i] = value; + } + + private: + umpire::ResourceManager& m_resource_manager{umpire::ResourceManager::getInstance()}; + umpire::Allocator m_allocator{m_resource_manager.getAllocator("DEVICE")}; + std::size_t m_size{0}; + ElementT* m_data{nullptr}; + ExecutionContext m_touch{NONE}; + }; // class PinnedMemoryManager +} // namespace expt +} // namespace chai + +#endif // CHAI_PINNED_MEMORY_MANAGER_HPP diff --git a/src/chai/expt/UnifiedArrayManager.hpp b/src/chai/expt/UnifiedArrayManager.hpp new file mode 100644 index 00000000..93e168df --- /dev/null +++ b/src/chai/expt/UnifiedArrayManager.hpp @@ -0,0 +1,169 @@ +#ifndef CHAI_UNIFIED_ARRAY_MANAGER_HPP +#define CHAI_UNIFIED_ARRAY_MANAGER_HPP + +#include "chai/expt/ArrayManager.hpp" +#include "chai/expt/ExecutionContext.hpp" +#include "umpire/ResourceManager.hpp" + +namespace chai { +namespace expt { + class UnifiedArrayManager : public ArrayManager { + public: + UnifiedArrayManager() = default; + + explicit UnifiedArrayManager(const umpire::Allocator& allocator) : + m_allocator{allocator} + { + } + + UnifiedArrayManager(std::size_t size, const umpire::Allocator& allocator) : + m_allocator{allocator}, + m_size{size}, + m_data{m_allocator.allocate(m_size)} + { + } + + explicit UnifiedArrayManager(int allocatorID) : + m_allocator{m_resource_manager.getAllocator(allocatorID)} + { + } + + UnifiedArrayManager(std::size_t size, int allocatorID) : + m_allocator{m_resource_manager.getAllocator(allocatorID)}, + m_size{size}, + m_data{m_allocator.allocate(m_size)} + { + m_data = m_allocator.allocate(size); + } + + UnifiedArrayManager(const UnifiedArrayManager& other) : + m_size{other.m_size}, + m_allocator{other.m_allocator} + { + m_data = m_allocator.allocate(m_size * sizeof(T)); + m_execution_context_manager.setExecutionContext(ExecutionContext::DEVICE); + m_resource_manager.copy(other.m_data, m_data, m_size * sizeof(T)); + m_execution_context_manager.setExecutionContext(ExecutionContext::NONE); + m_modified = ExecutionContext::DEVICE; + } + + UnifiedArrayManager(UnifiedArrayManager&& other) : + m_data{other.m_data}, + m_size{other.m_size}, + m_modified{other.m_modified}, + m_allocator{other.m_allocator} + { + other.m_data = nullptr; + other.m_size = 0; + other.m_modified = NONE; + other.m_allocator = umpire::Allocator(); + } + + UnifiedArrayManager& operator=(const UnifiedArrayManager& other) { + if (&other != this) { // Prevent self-assignment + m_allocator.deallocate(m_data); + + m_size = other.m_size; + m_allocator = other.m_allocator; + m_data = m_allocator.allocate(m_size * sizeof(T)); + m_execution_context_manager.setExecutionContext(ExecutionContext::DEVICE); + m_resource_manager.copy(other.m_data, m_data, m_size * sizeof(T)); + m_execution_context_manager.setExecutionContext(ExecutionContext::NONE); + m_modified = ExecutionContext::DEVICE; + } + + return *this; + } + + UnifiedArrayManager& operator=(UnifiedArrayManager&& other) { + if (&other != this) { + m_allocator.deallocate(m_data); + + m_data = other.m_data; + m_size = other.m_size; + m_modified = other.m_modified; + m_allocator = other.m_allocator; + + other.m_data = nullptr; + other.m_size = 0; + other.m_modified = ExecutionContext::NONE; + other.m_allocator = umpire::Allocator(); + } + + return *this; + } + + /*! + * \brief Destructor. + */ + ~UnifiedArrayManager() { + m_allocator.deallocate(m_data); + } + + /*! + * \brief Get the number of elements. + */ + size_t size() const { + return m_size; + } + + void* data() { + ExecutionContext executionContext = m_execution_context_manager.getExecutionContext(); + + if (executionContext != m_modified) { + m_execution_context_manager.synchronize(m_modified); + m_modified = executionContext; + } + + return m_data; + } + + const T* data(ExecutionContext executionContext) const { + if (executionContext != m_modified) { + m_execution_context_manager.synchronize(m_modified); + m_modified = ExecutionContext::NONE; + } + + return m_data; + } + + T& get(ExecutionContext executionContext, size_t i) { + if (executionContext != m_modified) { + m_execution_context_manager.synchronize(m_modified); + m_modified = executionContext; + } + + return m_data[i]; + } + + const T& get(ExecutionContext executionContext, size_t i) { + if (executionContext != m_modified) { + m_execution_context_manager.synchronize(m_modified); + m_modified = ExecutionContext::NONE; + } + + return m_data[i]; + } + + ExecutionContext getModified() { + return m_modified; + } + + umpire::Allocator getAllocator() { + return m_allocator; + } + + private: + umpire::ResourceManager& m_resource_manager{umpire::ResourceManager::getInstance()}; + umpire::Allocator m_allocator{}; + T* m_data{nullptr}; + size_t m_size{0}; + ExecutionContext m_modified{ExecutionContext::NONE}; + + ExecutionContextManager& m_execution_context_manager{ExecutionContextManager::getInstance()}; + + }; // class UnifiedArrayManager +} // namespace expt +} // namespace chai + +#endif // CHAI_UNIFIED_ARRAY_MANAGER_HPP diff --git a/src/chai/expt/UnifiedMemoryArray.hpp b/src/chai/expt/UnifiedMemoryArray.hpp new file mode 100644 index 00000000..4d76a436 --- /dev/null +++ b/src/chai/expt/UnifiedMemoryArray.hpp @@ -0,0 +1,202 @@ +#ifndef CHAI_UNIFIED_MEMORY_ARRAY_HPP +#define CHAI_UNIFIED_MEMORY_ARRAY_HPP + +#include "chai/expt/ExecutionContext.hpp" +#include "umpire/ResourceManager.hpp" + +namespace chai { +namespace expt { + /*! + * \class UnifiedMemoryArray + * + * \brief A container for managing the lifetime and coherence of a + * unified memory array, meaning an array with a single address + * that is accessible from all processors/devices in a system. + * + * This container should be used in tandem with the ExecutionContextManager. + * Together, they provide a programming model where work (e.g. a kernel) + * is generally performed asynchronously, with synchronization occurring + * only as needed for coherence of the array. For example, if the array is + * written to in an asynchronize kernel on a GPU, then the GPU will be + * synchronized if the array needs to be accessed on the CPU. + * + * This model works well for APUs where the CPU and GPU have the same + * physical memory. It also works for pinned (i.e. page-locked) memory + * and in some cases for pageable memory, though no pre-fetching is + * performed. + * + * Example: + * + * \code + * // Create a UnifiedMemoryArray with size 100 and default allocator + * int size = 10000; + * UnifiedMemoryArray array(size); + * + * // Access elements on the device + * std::span device_view(array.data(ExecutionContext::DEVICE, array.size()); + * + * // Launch a kernel that modifies device_view. + * // Note that this example relies on c++20 and the ability to use constexpr + * // host code on the device. + * + * // Access elements on the host. This will synchronize the device. + * std::span host_view(array.data(ExecutionContext::HOST), array.size()); + * + * for (int i = 0; i < size; ++i) { + * std::cout << host_view[i] << "\n"; + * } + * + * // Access and modify individual elements in the container. + * // This should be used sparingly or it will tank performance. + * // Getting the last element after performing a scan is one use case. + * array.get(ExecutionContext::HOST, size - 1) = 10; + * \endcode + */ + template + class UnifiedMemoryArray + public: + UnifiedMemoryArray() = default; + + explicit UnifiedMemoryArray(const umpire::Allocator& allocator) : + m_allocator{allocator} + { + } + + UnifiedMemoryArray(std::size_t size, const umpire::Allocator& allocator) : + m_size{size}, + m_allocator{allocator} + { + m_data = m_allocator.allocate(m_size * sizeof(T)); + // TODO: Investigate if/when to do initialization + } + + explicit UnifiedMemoryArray(int allocatorID) : + m_allocator{umpire::ResourceManager::getInstance().getAllocator(allocatorID)} + { + } + + UnifiedMemoryArray(std::size_t size, int allocatorID) : + m_size{size}, + m_allocator{umpire::ResourceManager::getInstance().getAllocator(allocatorID)} + { + m_data = m_allocator.allocate(m_size * sizeof(T)); + // TODO: Investigate if/when to do initialization + } + + UnifiedMemoryArray(const UnifiedMemoryArray& other) : + m_size{other.m_size}, + m_allocator{other.m_allocator} + { + m_data = m_allocator.allocate(m_size * sizeof(T)); + ExecutionContextManager::getInstance().setExecutionContext(ExecutionContext::DEVICE); + umpire::ResourceManager::getInstance().copy(other.m_data, m_data, m_size * sizeof(T)); + ExecutionContextManager::getInstance().setExecutionContext(ExecutionContext::NONE); + m_last_execution_context = ExecutionContext::DEVICE; + } + + UnifiedMemoryArray(UnifiedMemoryArray&& other) : + m_data{other.m_data}, + m_size{other.m_size}, + m_last_execution_context{other.m_last_execution_context}, + m_allocator{other.m_allocator} + { + other.m_data = nullptr; + other.m_size = 0; + other.m_last_execution_context = NONE; + other.m_allocator = umpire::Allocator(); + } + + UnifiedMemoryArray& operator=(const UnifiedMemoryArray& other) { + if (&other != this) { // Prevent self-assignment + m_allocator.deallocate(m_data); + + m_size = other.m_size; + m_allocator = other.m_allocator; + m_data = m_allocator.allocate(m_size * sizeof(T)); + ExecutionContextManager::getInstance().setExecutionContext(ExecutionContext::DEVICE); + umpire::ResourceManager::getInstance().copy(other.m_data, m_data, m_size * sizeof(T)); + ExecutionContextManager::getInstance().setExecutionContext(ExecutionContext::NONE); + m_last_execution_context = ExecutionContext::DEVICE; + } + + return *this; + } + + UnifiedMemoryArray& operator=(UnifiedMemoryArray&& other) { + if (&other != this) { + m_allocator.deallocate(m_data); + + m_data = other.m_data; + m_size = other.m_size; + m_last_execution_context = other.m_last_execution_context; + m_allocator = other.m_allocator; + + other.m_data = nullptr; + other.m_size = 0; + other.m_last_execution_context = ExecutionContext::NONE; + other.m_allocator = umpire::Allocator(); + } + + return *this; + } + + /*! + * \brief Destructor. + */ + ~UnifiedMemoryArray() { + m_allocator.deallocate(m_data); + } + + /*! + * \brief Get the number of elements. + */ + size_t size() const { + return m_size; + } + + T* data(ExecutionContext executionContext) { + if (executionContext != m_last_execution_context) { + ExecutionContextManager::getInstance().synchronize(m_last_execution_context); + m_last_execution_context = executionContext; + } + + return m_data; + } + + const T* data(ExecutionContext executionContext) const { + if (executionContext != m_last_execution_context) { + ExecutionContextManager::getInstance().synchronize(m_last_execution_context); + m_last_execution_context = ExecutionContext::NONE; + } + + return m_data; + } + + T& get(ExecutionContext executionContext, size_t i) { + if (executionContext != m_last_execution_context) { + ExecutionContextManager::getInstance().synchronize(m_last_execution_context); + m_last_execution_context = executionContext; + } + + return m_data[i]; + } + + const T& get(ExecutionContext executionContext, size_t i) { + if (executionContext != m_last_execution_context) { + ExecutionContextManager::getInstance().synchronize(m_last_execution_context); + m_last_execution_context = ExecutionContext::NONE; + } + + return m_data[i]; + } + + private: + T* m_data{nullptr}; + size_t m_size{0}; + ExecutionContext m_last_execution_context{ExecutionContext::NONE}; + umpire::Allocator m_allocator{}; + }; // class UnifiedMemoryArray +} // namespace expt +} // namespace chai + +#endif // CHAI_UNIFIED_MEMORY_ARRAY_HPP diff --git a/src/chai/expt/UnifiedMemoryArrayPointer.hpp b/src/chai/expt/UnifiedMemoryArrayPointer.hpp new file mode 100644 index 00000000..f4ea3691 --- /dev/null +++ b/src/chai/expt/UnifiedMemoryArrayPointer.hpp @@ -0,0 +1,301 @@ +#ifndef CHAI_UNIFIED_MEMORY_POINTER_HPP +#define CHAI_UNIFIED_MEMORY_POINTER_HPP + +#include "chai/expt/UnifiedMemoryManager.hpp" + +#include +#include + +namespace chai { +namespace expt { + /*! + * \class UnifiedMemoryPointer + * + * \brief A data structure for managing the lifetime and coherence of a + * unified memory array, meaning an array with a single address + * that is accessible from all processors/devices in a system. + * + * This data structure is designed for use with the ExecutionContextManager. + * When the execution context is set and the copy constructor is triggered, + * the underlying data is made coherent in the current execution context. + * + * This lends itself particularly well to an elegant programming model + * where loop bodies are replaced with lambda expressions that capture + * variables in the current scope by copy. If the execution context has + * already been set, the lambda expression can be executed in that context. + * Otherwise, the execution context must be set and then a copy of the + * lambda expression (which triggers a copy of all the captured variables) + * can be executed in that context. This latter case is how CHAI works + * with RAJA via a RAJA plugin. + * + * This model works well for APUs where the CPU and GPU have the same + * physical memory. It also works for pinned (i.e. page-locked) memory. + * + * Example using RAJA and CHAI: + * + * \code + * #include "chai/UnifiedMemoryPointer.hpp" + * #include "RAJA/RAJA.hpp" + * + * constexpr int CUDA_BLOCK_SIZE = 256; + * constexpr int ASYNCHRONOUS = true; + * + * int size = 10000; + * chai::UnifiedMemoryPointer a(size); + * + * int offset = 42; + * + * // Both `a` and `offset` are captured by copy into the lambda expression. + * // The execution context is not set, so the copy constructor of `a` + * // results in a shallow copy with nothing done related to coherence + * // management. RAJA then calls the CHAI plugin, which sets the current + * // execution context. At that point, the lambda expression is then copied, + * // which again triggers the copy constructor of `a`. This time, the data + * // is made coherent on the device (which is essentially a no-op because it + * // has not been accessed in any other execution context yet). After this, + * // RAJA calls the CHAI plugin to reset the execution context and launches + * // a CUDA kernel that executes the lambda expression for each index in + * // [0, size). + * RAJA::forall>(RAJA::TypedRangeSegment(0, size), [=] __device__ (int i) { + * a[i] = offset + i; + * }); + * + * // The same process as described above happens, except that now `a` is + * // made coherent on the host by synchronizing the device and the lambda + * // expression is evaluated on the host. + * RAJA::forall(RAJA::TypedRangeSegment(0, size), [=] __device__ (int i) { + * a[i] = i - offset; + * }); + * \endcode + */ + template + class UnifiedMemoryPointer { + public: + /*! + * \brief Constructs an empty array. + */ + constexpr UnifiedMemoryPointer() noexcept = default; + + explicit UnifiedMemoryPointer(const umpire::Allocator& allocator) + : m_manager{new UnifiedMemoryManager(allocator)} + { + } + + UnifiedMemoryPointer(std::size_t size, const umpire::Allocator& allocator) + : m_size{size}, + m_manager{new UnifiedMemoryManager(size * sizeof(T), allocator)} + { + } + + explicit UnifiedMemoryPointer(int allocatorID) + : m_manager{new UnifiedMemoryManager(allocatorID)} + { + } + + UnifiedMemoryPointer(std::size_t size, int allocatorID) + : m_size{size}, + m_manager{new UnifiedMemoryManager(size * sizeof(T), allocatorID)} + { + } + + explicit UnifiedMemoryPointer(UnifiedMemoryManager* manager) + : m_manager{manager} + { + if (m_manager) + { + m_size = m_manager->size() / sizeof(T); + } + } + + /*! + * \brief Constructs a shallow copy of an array from another and makes + * the data coherent in the current execution space. + * + * \param other The other array. + * + * \note This is a shallow copy. + */ + CHAI_HOST_DEVICE UnifiedMemoryPointer(const UnifiedMemoryPointer& other) : + m_data{other.m_data}, + m_size{other.m_size} +#if !defined(CHAI_DEVICE_COMPILE) + , m_manager{other.m_manager} +#endif + { + update(); + } + + /*! + * \brief Sets the array manager for this UnifiedMemoryPointer. + * + * \param manager The new array manager to be set. + * + * \post The UnifiedMemoryPointer takes ownership of the new manager objet. + */ + void setManager(UnifiedMemoryManager* manager) + { + delete m_manager; + m_manager = manager; + + if (m_manager) + { + m_size = m_manager->size() / sizeof(T); + } + } + + /*! + * \brief Get the array manager associated with this UnifiedMemoryPointer. + * + * \return A pointer to the array manager. + */ + UnifiedMemoryManager* getManager() const + { + return m_manager; + } + + /*! + * \brief Resizes the array to the specified new size. + * + * \param newSize The new size to resize the array to. + * + * \note This method updates the size of the array and triggers a resize operation in the array manager if it exists. + * If no array manager is associated, an exception is thrown. + */ + void resize(std::size_t newSize) + { + if (m_manager) { + m_size = newSize; + m_manager->resize(newSize * sizeof(T)); + } + else { + m_manager = new UnifiedMemoryManager(newSize * sizeof(T)); + } + } + + /*! + * \brief Frees the resources associated with this array. + * + * \note Once free has been called, it is invalid to use any other copies + * of this array (since copies are shallow). + */ + void free() { + m_data = nullptr; + m_size = 0; + delete m_manager; + m_manager = nullptr; + } + + /*! + * \brief Get the number of elements in the array. + * + * \pre The copy constructor has been called with the execution space + * set to CPU or GPU (e.g. by the RAJA plugin). + */ + CHAI_HOST_DEVICE std::size_t size() const + { + return m_size; + } + + CHAI_HOST_DEVICE update() const + { +#if !defined(CHAI_DEVICE_COMPILE) + if (m_manager) { + m_data = static_cast(m_manager->data(!std::is_const::value)); + // m_size = m_manager->size() / sizeof(T); + } +#endif + } + + /*! + * \brief Get a pointer to the element data in the current execution context. + * + * \return A pointer to the element data in the current execution context. + */ + CHAI_HOST_DEVICE T* data() const + { + update(); + return m_data; + } + + /*! + * \brief Get a const pointer to the element data in the current execution context. + * + * \return A const pointer to the element data in the current execution context. + */ + CHAI_HOST_DEVICE const T* cdata() const { +#if !defined(CHAI_DEVICE_COMPILE) + if (m_manager) { + m_data = static_cast(m_manager->data(false)); + m_size = m_manager->size() / sizeof(T); + } +#endif + return m_data; + } + + /*! + * \brief Get the ith element in the array. + * + * \param i The index of the element to retrieve. + * + * \pre The copy constructor has been called with the execution space + * set to CPU or GPU (e.g. by the RAJA plugin). + */ + CHAI_HOST_DEVICE T& operator[](std::size_t i) const { + return m_data[i]; + } + + /*! + * \brief Get the value of the element at the specified index. + * + * \param i The index of the element to retrieve. + * + * \return The value of the element at the specified index. + * + * \throw std::runtime_exception if unable to retrieve the element. + */ + T get(std::size_t i) const { + if (m_manager) { + return m_manager->get(i); + } + else { + throw std::out_of_range(); + } + } + + /*! + * \brief Set a value at a specified index in the array. + * + * \param i The index where the value is to be set. + * \param value The value to set at the specified index. + * + * \throw std::runtime_exception if the array manager is not associated with the UnifiedMemoryPointer. + */ + void set(std::size_t i, const T& value) { + if (m_manager) { + m_manager->set(i, value); + } + else { + throw std::out_of_range(); + } + } + + private: + /*! + * The array that is coherent in the current execution space. + */ + T* m_data{nullptr}; + + /*! + * The number of elements in the array. + */ + std::size_t m_size{0}; + + /*! + * The array manager controls the coherence of the array. + */ + UnifiedMemoryManager* m_manager{nullptr}; + }; // class UnifiedMemoryPointer +} // namespace expt +} // namespace chai + +#endif // CHAI_UNIFIED_MEMORY_POINTER_HPP diff --git a/src/chai/expt/UnifiedMemoryManager.hpp b/src/chai/expt/UnifiedMemoryManager.hpp new file mode 100644 index 00000000..ef376194 --- /dev/null +++ b/src/chai/expt/UnifiedMemoryManager.hpp @@ -0,0 +1,168 @@ +#ifndef CHAI_UNIFIED_MEMORY_MANAGER_HPP +#define CHAI_UNIFIED_MEMORY_MANAGER_HPP + +#include "chai/expt/ExecutionContext.hpp" +#include "umpire/ResourceManager.hpp" + +namespace chai { +namespace expt { + class UnifiedMemoryManager + public: + UnifiedMemoryManager() = default; + + explicit UnifiedMemoryManager(const umpire::Allocator& allocator) : + m_allocator{allocator} + { + } + + UnifiedMemoryManager(std::size_t size, const umpire::Allocator& allocator) : + m_allocator{allocator}, + m_size{size}, + m_data{m_allocator.allocate(m_size)} + { + } + + explicit UnifiedMemoryManager(int allocatorID) : + m_allocator{m_resource_manager.getAllocator(allocatorID)} + { + } + + UnifiedMemoryManager(std::size_t size, int allocatorID) : + m_allocator{m_resource_manager.getAllocator(allocatorID)}, + m_size{size}, + m_data{m_allocator.allocate(m_size)} + { + m_data = m_allocator.allocate(size); + } + + UnifiedMemoryManager(const UnifiedMemoryManager& other) : + m_size{other.m_size}, + m_allocator{other.m_allocator} + { + m_data = m_allocator.allocate(m_size * sizeof(T)); + m_execution_context_manager.setExecutionContext(ExecutionContext::DEVICE); + m_resource_manager.copy(other.m_data, m_data, m_size * sizeof(T)); + m_execution_context_manager.setExecutionContext(ExecutionContext::NONE); + m_modified = ExecutionContext::DEVICE; + } + + UnifiedMemoryManager(UnifiedMemoryManager&& other) : + m_data{other.m_data}, + m_size{other.m_size}, + m_modified{other.m_modified}, + m_allocator{other.m_allocator} + { + other.m_data = nullptr; + other.m_size = 0; + other.m_modified = NONE; + other.m_allocator = umpire::Allocator(); + } + + UnifiedMemoryManager& operator=(const UnifiedMemoryManager& other) { + if (&other != this) { // Prevent self-assignment + m_allocator.deallocate(m_data); + + m_size = other.m_size; + m_allocator = other.m_allocator; + m_data = m_allocator.allocate(m_size * sizeof(T)); + m_execution_context_manager.setExecutionContext(ExecutionContext::DEVICE); + m_resource_manager.copy(other.m_data, m_data, m_size * sizeof(T)); + m_execution_context_manager.setExecutionContext(ExecutionContext::NONE); + m_modified = ExecutionContext::DEVICE; + } + + return *this; + } + + UnifiedMemoryManager& operator=(UnifiedMemoryManager&& other) { + if (&other != this) { + m_allocator.deallocate(m_data); + + m_data = other.m_data; + m_size = other.m_size; + m_modified = other.m_modified; + m_allocator = other.m_allocator; + + other.m_data = nullptr; + other.m_size = 0; + other.m_modified = ExecutionContext::NONE; + other.m_allocator = umpire::Allocator(); + } + + return *this; + } + + /*! + * \brief Destructor. + */ + ~UnifiedMemoryManager() { + m_allocator.deallocate(m_data); + } + + /*! + * \brief Get the number of elements. + */ + size_t size() const { + return m_size; + } + + void* data() { + ExecutionContext executionContext = m_execution_context_manager.getExecutionContext(); + + if (executionContext != m_modified) { + m_execution_context_manager.synchronize(m_modified); + m_modified = executionContext; + } + + return m_data; + } + + const T* data(ExecutionContext executionContext) const { + if (executionContext != m_modified) { + m_execution_context_manager.synchronize(m_modified); + m_modified = ExecutionContext::NONE; + } + + return m_data; + } + + T& get(ExecutionContext executionContext, size_t i) { + if (executionContext != m_modified) { + m_execution_context_manager.synchronize(m_modified); + m_modified = executionContext; + } + + return m_data[i]; + } + + const T& get(ExecutionContext executionContext, size_t i) { + if (executionContext != m_modified) { + m_execution_context_manager.synchronize(m_modified); + m_modified = ExecutionContext::NONE; + } + + return m_data[i]; + } + + ExecutionContext getModified() { + return m_modified; + } + + umpire::Allocator getAllocator() { + return m_allocator; + } + + private: + umpire::ResourceManager& m_resource_manager{umpire::ResourceManager::getInstance()}; + umpire::Allocator m_allocator{}; + T* m_data{nullptr}; + size_t m_size{0}; + ExecutionContext m_modified{ExecutionContext::NONE}; + + ExecutionContextManager& m_execution_context_manager{ExecutionContextManager::getInstance()}; + + }; // class UnifiedMemoryManager +} // namespace expt +} // namespace chai + +#endif // CHAI_UNIFIED_MEMORY_MANAGER_HPP diff --git a/src/chai/expt/span.hpp b/src/chai/expt/span.hpp new file mode 100644 index 00000000..cd9eec59 --- /dev/null +++ b/src/chai/expt/span.hpp @@ -0,0 +1,96 @@ +#pragma once + +#include +#include + +namespace custom { + +template +class span { +public: + using element_type = T; + using value_type = std::remove_cv_t; + using size_type = std::size_t; + using difference_type = std::ptrdiff_t; + using pointer = T*; + using const_pointer = const T*; + using reference = T&; + using const_reference = const T&; + using iterator = pointer; + using const_iterator = const_pointer; + + // Constructors + __host__ __device__ constexpr span() noexcept : data_(nullptr), size_(0) {} + + __host__ __device__ constexpr span(pointer ptr, size_type count) noexcept + : data_(ptr), size_(count) {} + + template + __host__ __device__ constexpr span(element_type (&arr)[N]) noexcept + : data_(arr), size_(N) {} + + // Element access + __host__ __device__ constexpr reference operator[](size_type idx) const noexcept { + return data_[idx]; + } + + __host__ __device__ constexpr reference front() const noexcept { + return data_[0]; + } + + __host__ __device__ constexpr reference back() const noexcept { + return data_[size_ - 1]; + } + + __host__ __device__ constexpr pointer data() const noexcept { + return data_; + } + + // Iterators + __host__ __device__ constexpr iterator begin() const noexcept { + return data_; + } + + __host__ __device__ constexpr iterator end() const noexcept { + return data_ + size_; + } + + // Capacity + __host__ __device__ constexpr bool empty() const noexcept { + return size_ == 0; + } + + __host__ __device__ constexpr size_type size() const noexcept { + return size_; + } + + __host__ __device__ constexpr size_type size_bytes() const noexcept { + return size_ * sizeof(element_type); + } + + // Subviews + __host__ __device__ constexpr span first(size_type count) const { + return {data_, count}; + } + + __host__ __device__ constexpr span last(size_type count) const { + return {data_ + (size_ - count), count}; + } + + __host__ __device__ constexpr span subspan(size_type offset, size_type count) const { + return {data_ + offset, count}; + } + +private: + pointer data_; + size_type size_; +}; + +// Deduction guides +template +span(T (&)[N]) -> span; + +template +span(T*, std::size_t) -> span; + +} // namespace custom \ No newline at end of file diff --git a/tests/expt/UnifiedMemoryManagerTests.cpp b/tests/expt/UnifiedMemoryManagerTests.cpp new file mode 100644 index 00000000..baa5c59e --- /dev/null +++ b/tests/expt/UnifiedMemoryManagerTests.cpp @@ -0,0 +1,224 @@ +#include "gtest/gtest.h" +#include "chai/expt/UnifiedMemoryManager.hpp" + +class UnifiedMemoryManagerTest : public ::testing::Test { +protected: + void SetUp() override { + // Get a basic allocator for testing + m_allocator = umpire::ResourceManager::getInstance().getAllocator("HOST"); + m_execution_context_manager = ExecutionContextManager::getInstance(); + } + + umpire::Allocator m_allocator; + ExecutionContextManager& m_execution_context_manager; +}; + +TEST_F(UnifiedMemoryManagerTest, DefaultConstructor) { + chai::expt::UnifiedMemoryManager manager; + + { + EXPECT_EQ(manager.size(), 0); + EXPECT_EQ(manager.data(), nullptr); + } + + { + chai::expt::ExecutionContextGuard executionContextGuard(ExecutionContext::NONE); + EXPECT_EQ(manager.size(), 0); + EXPECT_EQ(manager.data(), nullptr); + } + + { + chai::expt::ExecutionContextGuard executionContextGuard(ExecutionContext::HOST); + EXPECT_EQ(manager.size(), 0); + EXPECT_EQ(manager.data(), nullptr); + } + +#if defined(CHAI_ENABLE_DEVICE) + { + chai::expt::ExecutionContextGuard executionContextGuard(ExecutionContext::DEVICE); + EXPECT_EQ(manager.size(), 0); + EXPECT_EQ(manager.data(), nullptr); + } +#endif +} + +TEST_F(UnifiedMemoryManagerTest, AllocatorConstructor) { + chai::expt::UnifiedMemoryManager manager(m_allocator); + EXPECT_EQ(manager.size(), 0); + EXPECT_EQ(manager.data(chai::expt::ExecutionContext::HOST), nullptr); +} + +TEST_F(UnifiedMemoryManagerTest, SizeAndAllocatorConstructor) { + const size_t size = 10; + chai::expt::UnifiedMemoryManager manager(size, m_allocator); + EXPECT_EQ(manager.size(), size); + EXPECT_NE(manager.data(chai::expt::ExecutionContext::HOST), nullptr); +} + +TEST_F(UnifiedMemoryManagerTest, AllocatorIDConstructor) { + chai::expt::UnifiedMemoryManager manager(0); // 0 typically corresponds to HOST + EXPECT_EQ(manager.size(), 0); + EXPECT_EQ(manager.data(chai::expt::ExecutionContext::HOST), nullptr); +} + +TEST_F(UnifiedMemoryManagerTest, SizeAndAllocatorIDConstructor) { + const size_t size = 10; + chai::expt::UnifiedMemoryManager manager(size, 0); // 0 typically corresponds to HOST + EXPECT_EQ(manager.size(), size); + EXPECT_NE(manager.data(chai::expt::ExecutionContext::HOST), nullptr); +} + +TEST_F(UnifiedMemoryManagerTest, CopyConstructor) { + const size_t size = 10; + chai::expt::UnifiedMemoryManager original(size, m_allocator); + + // Initialize data + for (size_t i = 0; i < size; ++i) { + original.get(chai::expt::ExecutionContext::HOST, i) = static_cast(i); + } + + // Copy + chai::expt::UnifiedMemoryManager copy(original); + + // Verify + EXPECT_EQ(copy.size(), original.size()); + for (size_t i = 0; i < size; ++i) { + EXPECT_EQ(copy.get(chai::expt::ExecutionContext::HOST, i), + original.get(chai::expt::ExecutionContext::HOST, i)); + } +} + +TEST_F(UnifiedMemoryManagerTest, MoveConstructor) { + const size_t size = 10; + chai::expt::UnifiedMemoryManager original(size, m_allocator); + + // Initialize data + for (size_t i = 0; i < size; ++i) { + original.get(chai::expt::ExecutionContext::HOST, i) = static_cast(i); + } + + // Store data pointer for comparison + int* originalData = original.data(chai::expt::ExecutionContext::HOST); + + // Move + chai::expt::UnifiedMemoryManager moved(std::move(original)); + + // Verify + EXPECT_EQ(moved.size(), size); + EXPECT_EQ(moved.data(chai::expt::ExecutionContext::HOST), originalData); + EXPECT_EQ(original.size(), 0); + EXPECT_EQ(original.data(chai::expt::ExecutionContext::HOST), nullptr); +} + +TEST_F(UnifiedMemoryManagerTest, CopyAssignment) { + const size_t size = 10; + chai::expt::UnifiedMemoryManager original(size, m_allocator); + + // Initialize data + for (size_t i = 0; i < size; ++i) { + original.get(chai::expt::ExecutionContext::HOST, i) = static_cast(i); + } + + // Copy assignment + chai::expt::UnifiedMemoryManager copy; + copy = original; + + // Verify + EXPECT_EQ(copy.size(), original.size()); + for (size_t i = 0; i < size; ++i) { + EXPECT_EQ(copy.get(chai::expt::ExecutionContext::HOST, i), + original.get(chai::expt::ExecutionContext::HOST, i)); + } +} + +TEST_F(UnifiedMemoryManagerTest, MoveAssignment) { + const size_t size = 10; + chai::expt::UnifiedMemoryManager original(size, m_allocator); + + // Initialize data + for (size_t i = 0; i < size; ++i) { + original.get(chai::expt::ExecutionContext::HOST, i) = static_cast(i); + } + + // Store data pointer for comparison + int* originalData = original.data(chai::expt::ExecutionContext::HOST); + + // Move assignment + chai::expt::UnifiedMemoryManager moved; + moved = std::move(original); + + // Verify + EXPECT_EQ(moved.size(), size); + EXPECT_EQ(moved.data(chai::expt::ExecutionContext::HOST), originalData); + EXPECT_EQ(original.size(), 0); + EXPECT_EQ(original.data(chai::expt::ExecutionContext::HOST), nullptr); +} + +TEST_F(UnifiedMemoryManagerTest, DataAccess) { + const size_t size = 10; + chai::expt::UnifiedMemoryManager manager(size, m_allocator); + + // Initialize and verify data access + for (size_t i = 0; i < size; ++i) { + manager.get(chai::expt::ExecutionContext::HOST, i) = static_cast(i); + } + + // Verify using get() + for (size_t i = 0; i < size; ++i) { + EXPECT_EQ(manager.get(chai::expt::ExecutionContext::HOST, i), static_cast(i)); + } + + // Verify using data() + int* data = manager.data(chai::expt::ExecutionContext::HOST); + for (size_t i = 0; i < size; ++i) { + EXPECT_EQ(data[i], static_cast(i)); + } +} + +TEST_F(UnifiedMemoryManagerTest, ConstDataAccess) { + const size_t size = 10; + chai::expt::UnifiedMemoryManager manager(size, m_allocator); + + // Initialize data + for (size_t i = 0; i < size; ++i) { + manager.get(chai::expt::ExecutionContext::HOST, i) = static_cast(i); + } + + // Create const reference and verify data access + const chai::expt::UnifiedMemoryManager& constManager = manager; + + // Verify using get() + for (size_t i = 0; i < size; ++i) { + EXPECT_EQ(constManager.get(chai::expt::ExecutionContext::HOST, i), static_cast(i)); + } + + // Verify using data() + const int* constData = constManager.data(chai::expt::ExecutionContext::HOST); + for (size_t i = 0; i < size; ++i) { + EXPECT_EQ(constData[i], static_cast(i)); + } +} + +TEST_F(UnifiedMemoryManagerTest, ExecutionContextSwitching) { + // Note: This test assumes a system with both HOST and DEVICE execution contexts + // For systems without a device (e.g., GPU), this test may need to be modified + + const size_t size = 10; + chai::expt::UnifiedMemoryManager manager(size, m_allocator); + + // Initialize data on HOST + for (size_t i = 0; i < size; ++i) { + manager.get(chai::expt::ExecutionContext::HOST, i) = static_cast(i); + } + + // Access data on DEVICE (should trigger synchronization) + int* deviceData = manager.data(chai::expt::ExecutionContext::DEVICE); + + // Access data back on HOST (should trigger synchronization again) + int* hostData = manager.data(chai::expt::ExecutionContext::HOST); + + // Verify data is still correct + for (size_t i = 0; i < size; ++i) { + EXPECT_EQ(hostData[i], static_cast(i)); + } +} \ No newline at end of file diff --git a/tests/integration/CMakeLists.txt b/tests/integration/CMakeLists.txt index ca5dbb88..f861606b 100644 --- a/tests/integration/CMakeLists.txt +++ b/tests/integration/CMakeLists.txt @@ -25,6 +25,22 @@ blt_add_test( NAME managed_array_test COMMAND managed_array_tests) +if (FALSE) +# Array tests +blt_add_executable( + NAME TestArray + SOURCES TestArray.cpp + DEPENDS_ON ${chai_integration_test_depends}) + +target_include_directories( + TestArray + PUBLIC ${PROJECT_BINARY_DIR}/include) + +blt_add_test( + NAME TestArray + COMMAND TestArray) +endif() + if (CHAI_ENABLE_MANAGED_PTR) blt_add_executable( NAME managed_ptr_tests diff --git a/tests/integration/expt/TestArray.cpp b/tests/integration/expt/TestArray.cpp new file mode 100644 index 00000000..70c392ba --- /dev/null +++ b/tests/integration/expt/TestArray.cpp @@ -0,0 +1,73 @@ +////////////////////////////////////////////////////////////////////////////// +// Copyright (c) 2016-25, Lawrence Livermore National Security, LLC and CHAI +// project contributors. See the CHAI LICENSE file for details. +// +// SPDX-License-Identifier: BSD-3-Clause +////////////////////////////////////////////////////////////////////////////// +#include "gtest/gtest.h" +#define GPU_TEST(X, Y) \ + static void gpu_test_##X##Y(); \ + TEST(X, Y) { gpu_test_##X##Y(); } \ + static void gpu_test_##X##Y() + +#ifdef NDEBUG + +#ifdef CHAI_ENABLE_CUDA +#define device_assert(EXP) if( !(EXP) ) asm ("trap;") +#else +#define device_assert(EXP) if( !(EXP) ) asm ("s_trap 1;") +#endif + +#else +#define device_assert(EXP) assert(EXP) +#endif + +#ifdef CHAI_DISABLE_RM +#define assert_empty_map(IGNORED) +#else +#define assert_empty_map(IGNORED) ASSERT_EQ(chai::ArrayManager::getInstance()->getPointerMap().size(),0) +#endif + + +#include "chai/config.hpp" + +#include "../src/util/forall.hpp" + +#include "chai/ManagedArray.hpp" + +#include "umpire/ResourceManager.hpp" + + +TEST(Array, HostManager) +{ + chai::Array a = chai::makeArray(10, allocator); + + forall(sequential(), 0, 10, [=](int i) { array[i] = i; }); + + forall(sequential(), 0, 10, [=](int i) { ASSERT_EQ(array[i], i); }); + + array.free(); + + assert_empty_map(true); +} + + +GPU_TEST(ManagedArray, PickandSetDeviceToDeviceUM) +{ + chai::ManagedArray array1(10, chai::UM); + chai::ManagedArray array2(10, chai::UM); + + forall(gpu(), 0, 10, [=] __device__(int i) { array1[i] = i; }); + + forall(gpu(), 0, 10, [=] __device__(int i) { + int temp = array1.pick(i); + array2.set(i, temp); + }); + + forall(sequential(), 0, 10, [=](int i) { ASSERT_EQ(array2[i], i); }); + + array1.free(); + array2.free(); + assert_empty_map(true); +} + diff --git a/tests/unit/CMakeLists.txt b/tests/unit/CMakeLists.txt index 1831dfe9..14b8f161 100644 --- a/tests/unit/CMakeLists.txt +++ b/tests/unit/CMakeLists.txt @@ -50,3 +50,15 @@ if (CHAI_ENABLE_MANAGED_PTR) NAME managed_ptr_unit_test COMMAND managed_ptr_unit_tests) endif () + +if (CHAI_ENABLE_EXPERIMENTAL) + blt_add_executable(NAME ExecutionContextManagerTests + SOURCES expt/ExecutionContextManagerTests.cpp + DEPENDS_ON chai gtest) + + target_include_directories(ExecutionContextManagerTests + PUBLIC ${PROJECT_BINARY_DIR}/include) + + blt_add_test(NAME ExecutionContextManagerTests + COMMAND ExecutionContextManagerTests) +endif () diff --git a/tests/unit/expt/CopyHidingArrayManagerTests.cpp b/tests/unit/expt/CopyHidingArrayManagerTests.cpp new file mode 100644 index 00000000..c8702cb5 --- /dev/null +++ b/tests/unit/expt/CopyHidingArrayManagerTests.cpp @@ -0,0 +1,31 @@ +////////////////////////////////////////////////////////////////////////////// +// Copyright (c) 2016-25, Lawrence Livermore National Security, LLC and CHAI +// project contributors. See the CHAI LICENSE file for details. +// +// SPDX-License-Identifier: BSD-3-Clause +////////////////////////////////////////////////////////////////////////////// + +#include "chai/expt/CopyHidingArrayManager.hpp" +#include "gtest/gtest.h" + +TEST(CopyHidingArrayManager, DefaultConstructor) { + chai::expt::CopyingHidingArrayManager arrayManager{}; + EXPECT_EQ(arrayManager.size(), 0); + EXPECT_EQ(arrayManager.data(chai::expt::ExecutionContext::NONE), nullptr); + EXPECT_EQ(arrayManager.data(chai::expt::ExecutionContext::HOST), nullptr); + EXPECT_EQ(arrayManager.data(chai::expt::ExecutionContext::DEVICE), nullptr); +} + +// Test that the default execution context is NONE +TEST(ExecutionContextManager, DefaultExecutionContext) { + chai::expt::ExecutionContextManager& executionContextManager = chai::expt::ExecutionContextManager::getInstance(); + EXPECT_EQ(executionContextManager.getExecutionContext(), chai::expt::ExecutionContext::NONE); +} + +// Test setting and getting the execution context +TEST(ExecutionContextManager, ExecutionContext) { + chai::expt::ExecutionContextManager& executionContextManager = chai::expt::ExecutionContextManager::getInstance(); + chai::expt::ExecutionContext executionContext = chai::expt::ExecutionContext::HOST; + executionContextManager.setExecutionContext(executionContext); + EXPECT_EQ(executionContextManager.getExecutionContext(), executionContext); +} \ No newline at end of file diff --git a/tests/unit/expt/ExecutionContextManagerTests.cpp b/tests/unit/expt/ExecutionContextManagerTests.cpp new file mode 100644 index 00000000..35802a66 --- /dev/null +++ b/tests/unit/expt/ExecutionContextManagerTests.cpp @@ -0,0 +1,30 @@ +////////////////////////////////////////////////////////////////////////////// +// Copyright (c) 2016-25, Lawrence Livermore National Security, LLC and CHAI +// project contributors. See the CHAI LICENSE file for details. +// +// SPDX-License-Identifier: BSD-3-Clause +////////////////////////////////////////////////////////////////////////////// + +#include "chai/expt/ExecutionContextManager.hpp" +#include "gtest/gtest.h" + +// Test that getInstance returns the same object at the same place in memory +TEST(ExecutionContextManager, SingletonInstance) { + chai::expt::ExecutionContextManager& executionContextManager1 = chai::expt::ExecutionContextManager::getInstance(); + chai::expt::ExecutionContextManager& executionContextManager2 = chai::expt::ExecutionContextManager::getInstance(); + EXPECT_EQ(&executionContextManager1, &executionContextManager2); +} + +// Test that the default execution context is NONE +TEST(ExecutionContextManager, DefaultExecutionContext) { + chai::expt::ExecutionContextManager& executionContextManager = chai::expt::ExecutionContextManager::getInstance(); + EXPECT_EQ(executionContextManager.getExecutionContext(), chai::expt::ExecutionContext::NONE); +} + +// Test setting and getting the execution context +TEST(ExecutionContextManager, ExecutionContext) { + chai::expt::ExecutionContextManager& executionContextManager = chai::expt::ExecutionContextManager::getInstance(); + chai::expt::ExecutionContext executionContext = chai::expt::ExecutionContext::HOST; + executionContextManager.setExecutionContext(executionContext); + EXPECT_EQ(executionContextManager.getExecutionContext(), executionContext); +} \ No newline at end of file diff --git a/tests/unit/expt/TestCopyHidingArray.cpp b/tests/unit/expt/TestCopyHidingArray.cpp new file mode 100644 index 00000000..04005e51 --- /dev/null +++ b/tests/unit/expt/TestCopyHidingArray.cpp @@ -0,0 +1,181 @@ +#include "gtest/gtest.h" +#include "chai/expt/CopyHidingArray.hpp" +#include + +namespace chai { +namespace expt { + +// Fixture class for CopyHidingArray tests +class CopyHidingArrayTest : public ::testing::Test { +protected: + void SetUp() override { + // Get ResourceManager instance + m_resource_manager = &umpire::ResourceManager::getInstance(); + + // Get default allocators + m_cpu_allocator = m_resource_manager->getAllocator("HOST"); + m_gpu_allocator = m_resource_manager->getAllocator("DEVICE"); + } + + umpire::ResourceManager* m_resource_manager; + umpire::Allocator m_cpu_allocator; + umpire::Allocator m_gpu_allocator; + + // Helper to fill array with test data + void fillWithTestData(CopyHidingArray& arr, int start_val = 1) { + int* data = arr.data(ExecutionContext::CPU); + for (size_t i = 0; i < arr.size(); ++i) { + data[i] = start_val + i; + } + } + + // Helper to verify array contents + void verifyArrayContents(const CopyHidingArray& arr, int start_val = 1) { + const int* data = arr.data(ExecutionContext::CPU); + for (size_t i = 0; i < arr.size(); ++i) { + EXPECT_EQ(data[i], start_val + i); + } + } +}; + +// Test default constructor +TEST_F(CopyHidingArrayTest, DefaultConstructor) { + CopyHidingArray arr; + EXPECT_EQ(arr.size(), 0); + EXPECT_EQ(arr.data(ExecutionContext::NONE), nullptr); +} + +// Test constructor with allocators +TEST_F(CopyHidingArrayTest, AllocatorConstructor) { + CopyHidingArray arr(m_cpu_allocator, m_gpu_allocator); + EXPECT_EQ(arr.size(), 0); +} + +// Test size constructor +TEST_F(CopyHidingArrayTest, SizeConstructor) { + const size_t size = 100; + CopyHidingArray arr(size); + EXPECT_EQ(arr.size(), size); +} + +// Test size and allocator constructor +TEST_F(CopyHidingArrayTest, SizeAndAllocatorConstructor) { + const size_t size = 100; + CopyHidingArray arr(size, m_cpu_allocator, m_gpu_allocator); + EXPECT_EQ(arr.size(), size); +} + +// Test copy constructor +TEST_F(CopyHidingArrayTest, CopyConstructor) { + const size_t size = 100; + CopyHidingArray arr1(size); + fillWithTestData(arr1); + + CopyHidingArray arr2(arr1); + EXPECT_EQ(arr2.size(), size); + verifyArrayContents(arr2); +} + +// Test move constructor +TEST_F(CopyHidingArrayTest, MoveConstructor) { + const size_t size = 100; + CopyHidingArray arr1(size); + fillWithTestData(arr1); + + CopyHidingArray arr2(std::move(arr1)); + EXPECT_EQ(arr2.size(), size); + EXPECT_EQ(arr1.size(), 0); + verifyArrayContents(arr2); +} + +// Test copy assignment +TEST_F(CopyHidingArrayTest, CopyAssignment) { + const size_t size = 100; + CopyHidingArray arr1(size); + fillWithTestData(arr1); + + CopyHidingArray arr2; + arr2 = arr1; + EXPECT_EQ(arr2.size(), size); + verifyArrayContents(arr2); +} + +// Test move assignment +TEST_F(CopyHidingArrayTest, MoveAssignment) { + const size_t size = 100; + CopyHidingArray arr1(size); + fillWithTestData(arr1); + + CopyHidingArray arr2; + arr2 = std::move(arr1); + EXPECT_EQ(arr2.size(), size); + EXPECT_EQ(arr1.size(), 0); + verifyArrayContents(arr2); +} + +// Test data access and coherence +TEST_F(CopyHidingArrayTest, DataCoherence) { + const size_t size = 100; + CopyHidingArray arr(size); + + // Fill on CPU + int* cpu_data = arr.data(ExecutionContext::CPU); + for (size_t i = 0; i < size; ++i) { + cpu_data[i] = i + 1; + } + + // Access on GPU (this will cause a copy) + int* gpu_data = arr.data(ExecutionContext::GPU); + // Here we would run a GPU kernel, but for testing we'll just verify the copy happens + + // Modify on GPU (simulation for test) + // In a real test, this would be done via a GPU kernel + for (size_t i = 0; i < size; ++i) { + gpu_data[i] *= 2; + } + + // Access back on CPU (should trigger a copy back) + cpu_data = arr.data(ExecutionContext::CPU); + + // Verify data was copied back correctly + for (size_t i = 0; i < size; ++i) { + EXPECT_EQ(cpu_data[i], (i + 1) * 2); + } +} + +// Test resize functionality +TEST_F(CopyHidingArrayTest, Resize) { + const size_t initial_size = 50; + const size_t new_size = 100; + + CopyHidingArray arr(initial_size); + fillWithTestData(arr); + + EXPECT_EQ(arr.size(), initial_size); + + arr.resize(new_size); + EXPECT_EQ(arr.size(), new_size); + + // First initial_size elements should still have their values + int* data = arr.data(ExecutionContext::CPU); + for (size_t i = 0; i < initial_size; ++i) { + EXPECT_EQ(data[i], i + 1); + } +} + +// Test const data access +TEST_F(CopyHidingArrayTest, ConstDataAccess) { + const size_t size = 100; + CopyHidingArray arr(size); + fillWithTestData(arr); + + const CopyHidingArray& const_arr = arr; + const int* const_data = const_arr.data(ExecutionContext::CPU); + + for (size_t i = 0; i < size; ++i) { + EXPECT_EQ(const_data[i], i + 1); + } +} + +} // namespace expt +} // namespace chai \ No newline at end of file diff --git a/tests/unit/expt/TestCopyHidingArrayManager.cpp b/tests/unit/expt/TestCopyHidingArrayManager.cpp new file mode 100644 index 00000000..f4217177 --- /dev/null +++ b/tests/unit/expt/TestCopyHidingArrayManager.cpp @@ -0,0 +1,11 @@ +////////////////////////////////////////////////////////////////////////////// +// Copyright (c) 2016-25, Lawrence Livermore National Security, LLC and CHAI +// project contributors. See the CHAI LICENSE file for details. +// +// SPDX-License-Identifier: BSD-3-Clause +////////////////////////////////////////////////////////////////////////////// +#include "gtest/gtest.h" + +#include "chai/config.hpp" +#include "chai/expt/CopyHidingManager.hpp" +#include "umpire/ResourceManager.hpp" \ No newline at end of file diff --git a/tests/unit/expt/TestCopyHidingManager.cpp b/tests/unit/expt/TestCopyHidingManager.cpp new file mode 100644 index 00000000..8c5c9e7f --- /dev/null +++ b/tests/unit/expt/TestCopyHidingManager.cpp @@ -0,0 +1,88 @@ +////////////////////////////////////////////////////////////////////////////// +// Copyright (c) 2016-25, Lawrence Livermore National Security, LLC and CHAI +// project contributors. See the CHAI LICENSE file for details. +// +// SPDX-License-Identifier: BSD-3-Clause +////////////////////////////////////////////////////////////////////////////// +#include "gtest/gtest.h" + +#include "chai/config.hpp" +#include "chai/expt/CopyHidingManager.hpp" +#include "umpire/ResourceManager.hpp" + +/*! + * CopyHidingManager has many states and transitions. A test fixture is created + * for each state, and a test case is created for each transition. + */ + +/*! + * \class CopyHidingManager_StateBothUnallocated_Test + * + * \brief Test fixture for the state where both host and device data are + * unallocated (and untouched) + */ +class CopyHidingManager_StateBothUnallocated_Test : public testing::Test { + protected: + int m_host_allocator_id{umpire::ResourceManager::getInstance().getAllocator("HOST").getId()}; + int m_device_allocator_id{umpire::ResourceManager::getInstance().getAllocator("DEVICE").getId()}; + std::size_t m_size{100}; + chai::expt::CopyHidingManager m_manager{m_host_allocator_id, + m_device_allocator_id, + m_size}; +}; + +TEST_F(CopyHidingManager_StateBothUnallocated_Test, Size) +{ + EXPECT_EQ(m_size, m_manager.size()); +} + +TEST_F(CopyHidingManager_StateBothUnallocated_Test, HostAllocatorID) +{ + EXPECT_EQ(m_host_allocator_id, m_manager.getHostAllocatorID()); +} + +TEST_F(CopyHidingManager_StateBothUnallocated_Test, DeviceAllocatorID) +{ + EXPECT_EQ(m_device_allocator_id, m_manager.getDeviceAllocatorID()); +} + +TEST_F(CopyHidingManager_StateBothUnallocated_Test, Touch) +{ + EXPECT_EQ(chai::expt::ExecutionContext::NONE, m_manager.getTouch()); +} + +TEST_F(CopyHidingManager_StateBothUnallocated_Test, Data_None) +{ + EXPECT_EQ(nullptr, m_manager.data(chai::expt::ExecutionContext::NONE, false)); + EXPECT_EQ(chai::expt::ExecutionContext::NONE, m_manager.getTouch()); +} + +TEST_F(CopyHidingManager_StateBothUnallocated_Test, Data_None_Touch) +{ + EXPECT_EQ(nullptr, m_manager.data(chai::expt::ExecutionContext::NONE, true)); + EXPECT_EQ(chai::expt::ExecutionContext::NONE, m_manager.getTouch()); +} + +TEST_F(CopyHidingManager_StateBothUnallocated_Test, Data_Host) +{ + EXPECT_NE(nullptr, m_manager.data(chai::expt::ExecutionContext::HOST, false)); + EXPECT_EQ(chai::expt::ExecutionContext::NONE, m_manager.getTouch()); +} + +TEST_F(CopyHidingManager_StateBothUnallocated_Test, Data_Host_Touch) +{ + EXPECT_NE(nullptr, m_manager.data(chai::expt::ExecutionContext::HOST, true)); + EXPECT_EQ(chai::expt::ExecutionContext::HOST, m_manager.getTouch()); +} + +TEST_F(CopyHidingManager_StateBothUnallocated_Test, Data_Device) +{ + EXPECT_NE(nullptr, m_manager.data(chai::expt::ExecutionContext::DEVICE, false)); + EXPECT_EQ(chai::expt::ExecutionContext::NONE, m_manager.getTouch()); +} + +TEST_F(CopyHidingManager_StateBothUnallocated_Test, Data_Device_Touch) +{ + EXPECT_NE(nullptr, m_manager.data(chai::expt::ExecutionContext::DEVICE, false)); + EXPECT_EQ(chai::expt::ExecutionContext::DEVICE, m_manager.getTouch()); +} diff --git a/tests/unit/expt/TestHostManager.cpp b/tests/unit/expt/TestHostManager.cpp new file mode 100644 index 00000000..ede568b8 --- /dev/null +++ b/tests/unit/expt/TestHostManager.cpp @@ -0,0 +1,46 @@ +////////////////////////////////////////////////////////////////////////////// +// Copyright (c) 2016-25, Lawrence Livermore National Security, LLC and CHAI +// project contributors. See the CHAI LICENSE file for details. +// +// SPDX-License-Identifier: BSD-3-Clause +////////////////////////////////////////////////////////////////////////////// +#include "gtest/gtest.h" + +#include "chai/config.hpp" +#include "chai/expt/HostManager.hpp" +#include "umpire/ResourceManager.hpp" + +class HostManagerTest : public testing::Test { + protected: + int m_allocator_id{umpire::ResourceManager::getInstance().getAllocator("HOST").getId()}; + std::size_t m_size{100}; + chai::expt::HostManager m_manager{m_allocator_id, m_size}; +}; + + +TEST_F(HostManagerTest, AllocatorID) +{ + EXPECT_EQ(m_allocator_id, m_manager.getAllocatorID()); +} + +TEST_F(HostManagerTest, Size) +{ + EXPECT_EQ(m_size, m_manager.size()); +} + +TEST_F(HostManagerTest, DataExecutionContextNone) +{ + EXPECT_EQ(nullptr, m_manager.data(chai::expt::ExecutionContext::NONE, false)); +} + +TEST_F(HostManagerTest, DataExecutionContextHost) +{ + EXPECT_NE(nullptr, m_manager.data(chai::expt::ExecutionContext::HOST, false)); +} + +#if defined(CHAI_ENABLE_CUDA) || defined(CHAI_ENABLE_HIP) || defined(CHAI_ENABLE_GPU_SIMULATION_MODE) +TEST_F(HostManagerTest, DataExecutionContextDevice) +{ + EXPECT_EQ(nullptr, m_manager.data(chai::expt::ExecutionContext::DEVICE, false)); +} +#endif diff --git a/tests/unit/expt/TestPinnedManager.cpp b/tests/unit/expt/TestPinnedManager.cpp new file mode 100644 index 00000000..6cc9ee62 --- /dev/null +++ b/tests/unit/expt/TestPinnedManager.cpp @@ -0,0 +1,128 @@ +////////////////////////////////////////////////////////////////////////////// +// Copyright (c) 2016-25, Lawrence Livermore National Security, LLC and CHAI +// project contributors. See the CHAI LICENSE file for details. +// +// SPDX-License-Identifier: BSD-3-Clause +////////////////////////////////////////////////////////////////////////////// +#include "gtest/gtest.h" + +#include "chai/config.hpp" +#include "chai/expt/HostManager.hpp" +#include "umpire/ResourceManager.hpp" + +class HostManagerTest : public testing::Test { + protected: + int m_allocator_id{umpire::ResourceManager::getInstance().getAllocator("HOST").getId()}; + std::size_t m_size{100}; + chai::expt::HostManager m_manager{m_allocator_id, m_size}; +}; + + +TEST_F(HostManagerTest, AllocatorID) +{ + EXPECT_EQ(m_allocator_id, m_manager.getAllocatorID()); +} + +TEST_F(HostManagerTest, Size) +{ + EXPECT_EQ(m_size, m_manager.size()); +} + +TEST_F(HostManagerTest, DataExecutionContextNone) +{ + EXPECT_EQ(nullptr, m_manager.data(chai::expt::ExecutionContext::NONE, false)); +} + +TEST_F(HostManagerTest, DataExecutionContextHost) +{ + EXPECT_NE(nullptr, m_manager.data(chai::expt::ExecutionContext::HOST, false)); +} + +#if defined(CHAI_ENABLE_CUDA) || defined(CHAI_ENABLE_HIP) || defined(CHAI_ENABLE_GPU_SIMULATION_MODE) +TEST_F(HostManagerTest, DataExecutionContextDevice) +{ + EXPECT_EQ(nullptr, m_manager.data(chai::expt::ExecutionContext::DEVICE, false)); +} +#endif + +TEST(PinnedManagerTest, Container1) +{ + // Could have accessor parameter to control whether operator[] is defined. + chai::ManagedArray myArray(10); + + { + ManagedView myView(myArray); + + RAJA_LOOP(i, 0, 10) { + // Could accidentally do a deep copy of myArray if didn't cast to a view + myView[i]++; + } RAJA_LOOP_END + } + + ManagedArray myArray2 = myArray; // Deep copy +} + +TEST(PinnedManagerTest, Container2) +{ + ManagedArray myArray(10); + + RAJA_LOOP(i, 0, 10, myView = ManagedView(myArray)) { + // Could accidentally do a deep copy of myArray if didn't cast to a view + myView[i]++; + } RAJA_LOOP_END + + ManagedArray myArray2 = myArray; // Deep copy +} + + +TEST(PinnedManagerTest, UniquePtr1) +{ + ManagedArray myArray(10); + + { + ManagedView myView(myArray); + + // Can't accidentally copy myArray into loop + RAJA_LOOP(i, 0, 10) { + myView[i]++; + } RAJA_LOOP_END + } + + ManagedArray& myArray2 = myArray; // Can't do a deep copy +} + +TEST(PinnedManagerTest, UniqueArray) +{ + chai::UniqueArray myArray = chai::makeUnique(10); + + // Can't accidentally copy myArray into loop + RAJA_LOOP(i, 0, 10, myView = ManagedView(myArray)) { + myView[i]++; + } RAJA_LOOP_END + + // Can't do a deep copy + ManagedArray myArray2 = myArray; // Deep copy +} + +TEST(PinnedManagerTest, SharedArray) +{ + chai::SharedArray myArray = chai::makeShared(10); + + RAJA_LOOP(i, 0, 10) { + // Can be used directly in RAJA loop + myArray[i]++; + } RAJA_LOOP_END +} + +TEST(PinnedManagerTest, NonOwnedArray) +{ + chai::NonOwnedArray myArray = chai::makeNonOwned(10); + + RAJA_LOOP(i, 0, 10) { + // Can be used directly in RAJA loop + myArray[i]++; + } RAJA_LOOP_END + + // Have to explicitly manage lifetime + myArray.free(); +}