diff --git a/docs/sphinx/expt/design.rst b/docs/sphinx/expt/design.rst index 705068ee..ca00854b 100644 --- a/docs/sphinx/expt/design.rst +++ b/docs/sphinx/expt/design.rst @@ -177,6 +177,50 @@ counted since clean up cannot be triggered from the device. a[i] -= 1; // Use CHAI data structures in the DEVICE context... }); +---------------- +ManagedArrayView +---------------- + +This class provides a uniform interface for viewing different types of memory +across multiple backends. It has shallow-copy semantics, which allows it to be +passed by value to a CUDA or HIP kernel. When copy constructed, it queries the +array manager to update the cached size and pointer from the array manager. +Unlike ``ManagedArrayPointer``, this class is non-owning: it cannot allocate, +resize, or free the underlying data, and the referenced manager must outlive +the view. + +.. code-block:: cpp + + #include "chai/expt/ContextRAJAPlugin.hpp" + #include "chai/expt/ManagedArrayView.hpp" + #include "chai/expt/UnifiedArrayManager.hpp" + #include "RAJA/RAJA.hpp" + + static ::RAJA::util::PluginRegistry::add P( + "CHAIContextPlugin", + "Plugin that integrates CHAI context management with RAJA."); + + // It's recommended to use an alias so that it is easy to swap out the array manager. + template + using UnifiedArrayManager = ::chai::expt::UnifiedArrayManager; + + template + using ManagedArrayView = ::chai::expt::ManagedArrayView>; + + const std::size_t N = 1000000; + UnifiedArrayManager manager{N}; + ManagedArrayView a{manager}; + + ::RAJA::forall<::RAJA::seq_exec>(::RAJA::TypedRangeSegment(0, N), [=] (int i) { + a[i] = i; // Use CHAI data structures in the HOST context... + }); + + constexpr int BLOCK_SIZE = 256; + + ::RAJA::forall<::RAJA::cuda_exec_async>(::RAJA::TypedRangeSegment(0, N), [=] __device__ (int i) { + a[i] -= 1; // Use CHAI data structures in the DEVICE context... + }); + ---------------- HostArrayManager ---------------- @@ -268,4 +312,3 @@ of each element on the host (numeric types are initialized to zero). const int* p = a.data(false); // Read through p... } - diff --git a/src/chai/CMakeLists.txt b/src/chai/CMakeLists.txt index 6b76eff9..549bd86f 100644 --- a/src/chai/CMakeLists.txt +++ b/src/chai/CMakeLists.txt @@ -42,6 +42,7 @@ if(CHAI_ENABLE_EXPERIMENTAL) expt/HostSharedPointer.hpp expt/ManagedArrayPointer.hpp expt/ManagedArraySharedPointer.hpp + expt/ManagedArrayView.hpp expt/UnifiedArrayManager.hpp ManagedSharedPtr.hpp SharedPtrCounter.hpp diff --git a/src/chai/expt/ManagedArrayView.hpp b/src/chai/expt/ManagedArrayView.hpp new file mode 100644 index 00000000..bc1612b2 --- /dev/null +++ b/src/chai/expt/ManagedArrayView.hpp @@ -0,0 +1,187 @@ +////////////////////////////////////////////////////////////////////////////// +// Copyright (c) Lawrence Livermore National Security, LLC and other CHAI +// contributors. See the CHAI LICENSE and COPYRIGHT files for details. +// +// SPDX-License-Identifier: BSD-3-Clause +////////////////////////////////////////////////////////////////////////////// + +#ifndef CHAI_EXPT_MANAGED_ARRAY_VIEW_HPP +#define CHAI_EXPT_MANAGED_ARRAY_VIEW_HPP + +#include "chai/config.hpp" +#include "chai/ChaiMacros.hpp" + +#include +#include +namespace chai::expt +{ + /*! + * \brief This class provides a uniform interface for viewing different types + * of managed array memory across multiple backends. + * + * \details This class has shallow-copy semantics, which allows it to be + * passed by value to a CUDA or HIP kernel. When copy constructed, + * it queries the array manager to update the cached size and pointer + * from the array manager. This acts as a non-owning view and cannot + * allocate, free, or resize the underlying data. + * + * \tparam ElementType The type of elements contained in this array. + * \tparam ManagerType Manages the underlying memory. + */ + template + class ManagedArrayView + { + public: + /*! + * \brief Constructs a default ManagedArrayView. + * + * \details Creates a null pointer with size zero and no associated manager. + */ + ManagedArrayView() = default; + + /*! + * \brief Constructs a ManagedArrayView from an existing manager. + * + * \param manager An object that owns/manages the underlying array. + * + * \note The manager is not copied or owned. The caller must ensure the + * manager outlives the view. + */ + explicit ManagedArrayView(ManagerType& manager) + : m_data{manager.data()}, + m_size{manager.size()}, + m_manager{&manager} + { + } + + /*! + * \brief Copy-constructs a ManagedArrayView from another ManagedArrayView. + * + * \param other The ManagedArrayView to copy. + * + * \details Copies the cached pointer, size, and manager pointer. The + * cached pointer/size are synchronized by calling update(). + */ + CHAI_HOST_DEVICE ManagedArrayView(const ManagedArrayView& other) + : m_data{other.m_data}, + m_size{other.m_size}, + m_manager{other.m_manager} + { + update(); + } + + /*! + * \brief Converting copy-constructor from a non-const ManagedArrayView to a const ManagedArrayView. + * + * \tparam OtherElementType The source element type; must be non-const, and this + * ManagedArrayView's ElementType must be const-qualified version of it. + * + * \param other The source ManagedArrayView to copy from. + */ + template && + std::is_same_v>>> + CHAI_HOST_DEVICE ManagedArrayView(const ManagedArrayView& other) + : m_data{other.m_data}, + m_size{other.m_size}, + m_manager{other.m_manager} + { + } + + /*! + * \brief Copy-assigns from another ManagedArrayView. + * + * \param other The ManagedArrayView to copy from. + * + * \return Reference to this ManagedArrayView. + */ + ManagedArrayView& operator=(const ManagedArrayView& other) = default; + + /*! + * \brief Returns the number of elements in the underlying managed array. + * + * \return The number of elements. + */ + CHAI_HOST_DEVICE std::size_t size() const + { +#if !defined(CHAI_DEVICE_COMPILE) + if (m_manager) + { + m_size = m_manager->size(); + } +#endif + return m_size; + } + + /*! + * \brief Returns the cached pointer to the managed array's data. + * + * \return Pointer to the first element of the underlying managed array, or nullptr. + */ + CHAI_HOST_DEVICE ElementType* data() const + { +#if !defined(CHAI_DEVICE_COMPILE) + if (m_manager) + { + if (ElementType* data = static_cast(m_manager->data()); data) + { + m_data = data; + } + } +#endif + return m_data; + } + + /*! + * \brief Synchronizes the cached pointer and size from the manager. + */ + CHAI_HOST_DEVICE void update() const + { +#if !defined(CHAI_DEVICE_COMPILE) + if (m_manager) + { + if (ElementType* data = static_cast(m_manager->data()); data) + { + m_data = data; + } + + m_size = m_manager->size(); + } +#endif + } + + /*! + * \brief Unchecked element access. + * + * \param i Element index. + * + * \return Reference to element i in the cached data pointer. + */ + CHAI_HOST_DEVICE ElementType& operator[](std::size_t i) const + { + return m_data[i]; + } + + private: + /*! + * \brief Cached pointer to the managed array. + */ + mutable ElementType* m_data{nullptr}; + + /*! + * \brief Cached number of elements in the managed array. + */ + mutable std::size_t m_size{0}; + + /*! + * \brief Pointer to the manager that owns/manages the underlying array. + */ + ManagerType* m_manager{nullptr}; + + /// Needed for the converting constructor + template + friend class ManagedArrayView; + }; // class ManagedArrayView +} // namespace chai::expt + +#endif // CHAI_EXPT_MANAGED_ARRAY_VIEW_HPP diff --git a/tests/expt/CMakeLists.txt b/tests/expt/CMakeLists.txt index 9e6049cb..19a1e993 100644 --- a/tests/expt/CMakeLists.txt +++ b/tests/expt/CMakeLists.txt @@ -70,6 +70,17 @@ blt_add_test( NAME ManagedArrayPointerTests COMMAND ManagedArrayPointerTests) +blt_add_executable( + NAME ManagedArrayViewTests + SOURCES ManagedArrayViewTests.cpp + HEADERS ${chai_expt_test_headers} + INCLUDES ${PROJECT_BINARY_DIR}/include + DEPENDS_ON ${chai_expt_test_depends}) + +blt_add_test( + NAME ManagedArrayViewTests + COMMAND ManagedArrayViewTests) + blt_add_executable( NAME HostArrayPointerTests SOURCES HostArrayPointerTests.cpp @@ -81,6 +92,17 @@ blt_add_test( NAME HostArrayPointerTests COMMAND HostArrayPointerTests) +blt_add_executable( + NAME HostArrayViewTests + SOURCES HostArrayViewTests.cpp + HEADERS ${chai_expt_test_headers} + INCLUDES ${PROJECT_BINARY_DIR}/include + DEPENDS_ON ${chai_expt_test_depends}) + +blt_add_test( + NAME HostArrayViewTests + COMMAND HostArrayViewTests) + blt_add_executable( NAME HostArraySharedPointerTests SOURCES HostArraySharedPointerTests.cpp diff --git a/tests/expt/HostArrayViewTests.cpp b/tests/expt/HostArrayViewTests.cpp new file mode 100644 index 00000000..c7e77031 --- /dev/null +++ b/tests/expt/HostArrayViewTests.cpp @@ -0,0 +1,179 @@ +////////////////////////////////////////////////////////////////////////////// +// Copyright (c) Lawrence Livermore National Security, LLC and other CHAI +// contributors. See the CHAI LICENSE and COPYRIGHT files for details. +// +// SPDX-License-Identifier: BSD-3-Clause +////////////////////////////////////////////////////////////////////////////// + +#include "chai/expt/HostArrayManager.hpp" +#include "chai/expt/ManagedArrayView.hpp" +#include "gtest/gtest.h" + +#include + +template +using HostArrayManager = ::chai::expt::HostArrayManager; + +template +using HostArrayView = ::chai::expt::ManagedArrayView>; + +template +using ConstHostArrayView = ::chai::expt::ManagedArrayView>; + +TEST(HostArrayView, DefaultConstructor) { + HostArrayView a; + EXPECT_EQ(a.size(), 0); + EXPECT_EQ(a.data(), nullptr); +} + +TEST(HostArrayView, ManagerDefaultConstructor) { + HostArrayManager manager; + HostArrayView a{manager}; + EXPECT_EQ(a.size(), 0); + EXPECT_EQ(a.data(), nullptr); +} + +TEST(HostArrayView, ManagerSizeConstructor) { + const std::size_t N = 10; + HostArrayManager manager{N}; + HostArrayView a{manager}; + EXPECT_EQ(a.size(), N); + ASSERT_NE(a.data(), nullptr); + + for (std::size_t i = 0; i < N; ++i) + { + a[i] = static_cast(i); + } +} + +TEST(HostArrayView, CopyConstructor) { + const std::size_t N = 10; + HostArrayManager manager{N}; + HostArrayView a{manager}; + + for (std::size_t i = 0; i < N; ++i) + { + a[i] = static_cast(i); + } + + HostArrayView b(a); + + EXPECT_EQ(b.size(), a.size()); + EXPECT_EQ(b.data(), a.data()); + + for (std::size_t i = 0; i < N; ++i) + { + EXPECT_EQ(b[i], static_cast(i)); + } + + for (std::size_t i = 0; i < N; ++i) + { + b[i] = -static_cast(i); + EXPECT_EQ(a[i], -static_cast(i)); + } +} + +TEST(HostArrayView, ConvertingConstructor) { + const std::size_t N = 10; + HostArrayManager manager{N}; + HostArrayView a{manager}; + + for (std::size_t i = 0; i < N; ++i) + { + a[i] = static_cast(i); + } + + ConstHostArrayView b(a); + + EXPECT_EQ(b.size(), a.size()); + EXPECT_EQ(b.data(), a.data()); + + for (std::size_t i = 0; i < N; ++i) + { + EXPECT_EQ(b[i], static_cast(i)); + } +} + +TEST(HostArrayView, CopyAssignmentOperator) { + HostArrayView a; + + const std::size_t N = 10; + HostArrayManager manager{N}; + a = HostArrayView(manager); + + EXPECT_EQ(a.size(), N); + EXPECT_NE(a.data(), nullptr); +} + +TEST(HostArrayView, Data) { + const std::size_t N = 10; + HostArrayManager manager{N}; + HostArrayView a{manager}; + + for (std::size_t i = 0; i < N; ++i) + { + a[i] = static_cast(i); + } + + int* data = a.data(); + + for (std::size_t i = 0; i < N; ++i) + { + EXPECT_EQ(data[i], static_cast(i)); + } +} + +TEST(HostArrayView, Update) { + const std::size_t N = 10; + HostArrayManager manager{N}; + HostArrayView a{manager}; + a.update(); + + for (std::size_t i = 0; i < N; ++i) + { + a[i] = static_cast(i); + } + + int* data = a.data(); + + for (std::size_t i = 0; i < N; ++i) + { + EXPECT_EQ(data[i], static_cast(i)); + } +} + +TEST(HostArrayView, ReflectsManagerResizeAfterUpdate) { + HostArrayManager manager{4}; + HostArrayView a{manager}; + + EXPECT_EQ(a.size(), 4); + + manager.resize(7); + a.update(); + + EXPECT_EQ(a.size(), 7); + EXPECT_EQ(a.data(), manager.data()); +} + +TEST(HostArrayView, LambdaCapture) { + const std::size_t N = 10; + HostArrayManager manager{N}; + HostArrayView a{manager}; + + auto f = [=] (std::size_t i) + { + a[i] = static_cast(i); + }; + + for (std::size_t i = 0; i < N; ++i) + { + f(i); + } + + int* data = a.data(); + + for (std::size_t i = 0; i < N; ++i) + { + EXPECT_EQ(data[i], static_cast(i)); + } +} diff --git a/tests/expt/ManagedArrayViewTests.cpp b/tests/expt/ManagedArrayViewTests.cpp new file mode 100644 index 00000000..838c43f2 --- /dev/null +++ b/tests/expt/ManagedArrayViewTests.cpp @@ -0,0 +1,159 @@ +////////////////////////////////////////////////////////////////////////////// +// Copyright (c) Lawrence Livermore National Security, LLC and other CHAI +// contributors. See the CHAI LICENSE and COPYRIGHT files for details. +// +// SPDX-License-Identifier: BSD-3-Clause +////////////////////////////////////////////////////////////////////////////// + +#include "chai/expt/ManagedArrayView.hpp" +#include "gtest/gtest.h" + +#include +#include + +namespace { + /** + * Minimal "ManagerType" for exercising ManagedArrayView in unit tests. + * + * Requirements satisfied (as used by ManagedArrayView): + * - std::size_t size() const + * - ElementType* data() + * + * Owns storage on host via std::realloc. + */ + template + class TestArrayManager + { + public: + TestArrayManager() = default; + + explicit TestArrayManager(std::size_t size) + { + m_size = size; + m_data = static_cast(std::realloc(m_data, size*sizeof(ElementType))); + } + + ~TestArrayManager() + { + std::free(m_data); + } + + void resize(std::size_t size) + { + m_size = size; + m_data = static_cast(std::realloc(m_data, size*sizeof(ElementType))); + } + + std::size_t size() const + { + return m_size; + } + + ElementType* data() + { + return m_data; + } + + private: + std::size_t m_size{0}; + ElementType* m_data{nullptr}; + }; // class TestArrayManager +} // anonymous namespace + +template +using TestArrayView = ::chai::expt::ManagedArrayView>; + +template +using ConstTestArrayView = ::chai::expt::ManagedArrayView>; + +TEST(ManagedArrayView, DefaultConstructor) { + TestArrayView a; + EXPECT_EQ(a.size(), 0); + EXPECT_EQ(a.data(), nullptr); +} + +TEST(ManagedArrayView, ManagerConstructor) { + TestArrayManager manager; + TestArrayView a{manager}; + EXPECT_EQ(a.size(), 0); + EXPECT_EQ(a.data(), nullptr); +} + +TEST(ManagedArrayView, CopyConstructor) { + TestArrayManager manager{10}; + TestArrayView a{manager}; + TestArrayView b(a); + EXPECT_EQ(b.size(), a.size()); + EXPECT_EQ(b.data(), a.data()); +} + +TEST(ManagedArrayView, ConvertingConstructor) { + TestArrayManager manager{10}; + TestArrayView a{manager}; + ConstTestArrayView b(a); + EXPECT_EQ(b.size(), a.size()); + EXPECT_EQ(b.data(), a.data()); +} + +TEST(ManagedArrayView, CopyAssignmentOperator) { + TestArrayView a; + TestArrayManager manager{10}; + a = TestArrayView(manager); + EXPECT_EQ(a.size(), 10); + EXPECT_NE(a.data(), nullptr); +} + +TEST(ManagedArrayView, Data) { + const std::size_t n = 10; + TestArrayManager manager{n}; + TestArrayView a{manager}; + a.update(); + + for (std::size_t i = 0; i < n; ++i) + { + a[i] = static_cast(i); + } + + int* data = a.data(); + + for (std::size_t i = 0; i < n; ++i) + { + EXPECT_EQ(data[i], static_cast(i)); + } +} + +TEST(ManagedArrayView, LambdaCapture) { + const std::size_t n = 10; + TestArrayManager manager{n}; + TestArrayView a{manager}; + + auto f = [=] (std::size_t i) + { + a[i] = static_cast(i); + }; + + for (std::size_t i = 0; i < n; ++i) + { + f(i); + } + + int* data = a.data(); + + for (std::size_t i = 0; i < n; ++i) + { + EXPECT_EQ(data[i], static_cast(i)); + } +} + +TEST(ManagedArrayView, ReflectsManagerResizeAfterUpdate) { + TestArrayManager manager{4}; + TestArrayView a{manager}; + + EXPECT_EQ(a.size(), 4); + + manager.resize(7); + a.update(); + + EXPECT_EQ(a.size(), 7); + EXPECT_EQ(a.data(), manager.data()); +}