diff --git a/src/chai/CMakeLists.txt b/src/chai/CMakeLists.txt
index 2285c544..3ac7e7ed 100644
--- a/src/chai/CMakeLists.txt
+++ b/src/chai/CMakeLists.txt
@@ -19,6 +19,8 @@ set (chai_headers
   ManagedArray.hpp
   ManagedArray.inl
   managed_ptr.hpp
+  PinnedArray.hpp
+  PinnedPlugin.hpp
   PointerRecord.hpp
   Types.hpp)
 
diff --git a/src/chai/PinnedArray.hpp b/src/chai/PinnedArray.hpp
new file mode 100644
index 00000000..cb61e227
--- /dev/null
+++ b/src/chai/PinnedArray.hpp
@@ -0,0 +1,72 @@
+#if !defined(CARE_PINNED_ARRAY_HPP)
+#define CARE_PINNED_ARRAY_HPP
+
+namespace care {
+
+   template <class T>
+   class PinnedArray {
+      public:
+         PinnedArray() = default;
+
+         PinnedArray(const Pinned& other) :
+            m_size{other.m_size},
+            m_data{other.m_data},
+            m_control{other.m_control}
+         {
+#if !defined(__HIP_DEVICE_COMPILE__)
+            // TODO: Implement callback
+            PinnedPlugin::registerCallback();
+#endif
+         }
+
+         void resize(size_t count) {
+            if (count == 0) {
+               free();
+            }
+            else if (m_size == 0) {
+               cuda::resources::Hip resource = cuda::resources::Hip::get_default();
+               hipMallocAsync(&m_data, count * sizeof(T), resource.get_stream());
+               m_control = new ControlBlock();
+               m_control->m_last_resource = resource;
+               m_control->m_last_event = resource.get_event();
+            }
+            else if (m_size != count) {
+               T* newData = nullptr;
+               hipMallocAsync(&newData, count * sizeof(T), m_control->m_last_resource.get_stream());
+               const size_t min = m_size < count ? m_size : count;
+               hipMemcpyAsync(newData, m_data, min * sizeof(T), hipMemcpyDeviceToDevice, m_control->m_last_resource);
+               m_control->m_last_event = m_control->m_last_resource.get_event();
+               hipFreeAsync(m_data, m_control->m_last_resource);
+               m_data = newData;
+            }
+
+            m_size = count;
+         }
+
+         void free() {
+            if (m_control) {
+               hipFreeAsync(m_data, m_control->m_last_resource.get_stream());
+               delete m_control;
+               m_control = nullptr;
+               m_data = nullptr;
+               m_size = 0;
+            }
+
+         T& operator[](size_t i) const {
+            return m_data[i];
+         }
+
+      private:
+         struct ControlBlock {
+            camp::resources::Resource m_last_resource;
+            camp::resources::Event m_last_event;
+         };
+
+         size_t m_size = 0;
+         T* m_data = nullptr;
+         ControlBlock* m_control = nullptr;
+   };
+
+}  // namespace care
+
+#endif  // CARE_PINNED_ARRAY_HPP
diff --git a/src/chai/PinnedArray2.hpp b/src/chai/PinnedArray2.hpp
new file mode 100644
index 00000000..d12fa27f
--- /dev/null
+++ b/src/chai/PinnedArray2.hpp
@@ -0,0 +1,230 @@
+#ifndef CHAI_PINNED_ARRAY_HPP
+#define CHAI_PINNED_ARRAY_HPP
+
+namespace chai {
+   class MemoryManager {
+      virtual void update(camp::resources::Resource& resource) = 0;
+
+      virtual void update(camp::resources::Resource& resource,
+                          camp::resources::Event& event) = 0;
+   };
+
+   class MemoryManagerPlugin : RAJA::util::PluginStrategy {
+      public:
+         void preCapture(camp::resources::Resource& resource) override {
+            m_isCapturing = true;
+         }
+
+         void postCapture(camp::resources::Resource& resource) override {
+            m_isCapturing = false;
+
+            for (MemoryManager* manager : m_managers) {
+               manager->update(resource);
+            }
+         }
+
+         void postLaunch(camp::resources::Resource& resource) override {
+            camp::resources::Event event = resource.get_event();
+
+            for (MemoryManager* manager : m_managers) {
+               manager->update(resource, event);
+            }
+
+            m_managers.clear();
+         }
+
+         void registerMemoryManager(MemoryManager* manager) {
+            if (m_isCapturing) {
+               m_managers.push_back(manager);
+            }
+         }
+
+      private:
+         bool m_isCapturing{false};
+         std::vector<MemoryManager*> m_managers;
+   };
+
+   class HipMallocMemoryManager : public MemoryManager
+   {
+      public:
+         HipMallocMemoryManager(size_t size, camp::resources::Hip& resource)
+           : MemoryManager(),
+             m_size(size),
+             m_last_resource(resource)
+         {
+            hipMallocAsync(&m_data, size, resource.get_stream());
+            m_event = resource.get_event();
+         }
+
+         ~HipMallocMemoryManager()
+         {
+            hipFreeAsync(m_data, m_last_resource.get_stream());
+         }
+
+         size_t size() const {
+            return m_size;
+         }
+
+         void update(camp::resources::Resource& resource) override {
+            if (resource != m_last_resource) {
+               m_event.wait();
+               m_event = camp::resources::Event();
+            }
+         }
+
+         void update(camp::resources::Resource& resource,
+                     camp::resources::Event& event) override {
+            m_last_resource = resource;
+            m_last_event = event;
+         }
+
+         void* data(camp::resources::Resource resource) {
+            if (resource != m_last_resource) {
+               m_event.wait();
+            }
+
+            return m_data;
+         }
+
+         void reset(camp::resources::Resource resource) {
+            m_last_resource = resource;
+            m_event = resource.get_event();
+         }
+
+      private:
+         size_t m_size = 0;
+         void* m_data = nullptr;
+         camp::resources::Resource m_last_resource;
+         camp::resources::Event m_event;
+   };
+
+   template <class T>
+   class PinnedArray {
+      public:
+         static PinnedArray allocate(size_t count,
+                                     camp::resources::Hip resource) {
+            if (count == 0) {
+               return PinnedArray();
+            }
+            else {
+               T* data;
+               hipMallocAsync(&data, count * sizeof(T), resource.get_stream());
+               ControlBlock* control = new ControlBlock(count, data, resource);
+               return PinnedArray(count, data, control);
+            }
+         }
+
+         static PinnedArray allocate(size_t count) {
+            return allocate(count, camp::resources::Hip::get_default());
+         }
+
+         static void deallocate(PinnedArray& ptr,
+                                camp::resources::Hip resource) {
+            ControlBlock* control = ptr.get_control();
+
+            if (control) {
+               if (resource != control->m_last_resource) {
+                  control->m_last_event.wait();
+               }
+
+               hipFreeAsync(control->m_data, resource.get_stream());
+               delete m_control;
+               ptr.reset();
+            }
+         }
+
+         PinnedArray() = default;
+
+         PinnedArray(size_t count,
+                     T* data,
+                     ControlBlock* control) :
+            m_size{count},
+            m_data{data},
+            m_control{control}
+         {
+         }
+
+         CHAI_HOST_DEVICE PinnedArray(const Pinned& other) :
+            m_size{other.m_size},
+            m_data{other.m_data},
+            m_control{other.m_control}
+         {
+#if !defined(__HIP_DEVICE_COMPILE__)
+            // TODO: Implement callback
+            PinnedPlugin::registerCallback();
+            m_data = static_cast<T*>(m_manager->data<T>());
+#endif
+         }
+
+         ControlBlock* get_control() {
+            return m_control;
+         }
+
+         void reset() {
+            m_size = 0;
+            m_data = nullptr;
+            m_control = nullptr;
+         }
+
+         void resize(size_t count) {
+            if (count == 0) {
+               free();
+            }
+            else if (m_size == 0) {
+               camp::resources::Hip resource = camp::resources::Hip::get_default();
+               hipMallocAsync(&m_data, count * sizeof(T), resource.get_stream());
+               m_control = new ControlBlock();
+               m_control->m_last_resource = resource;
+               m_control->m_last_event = resource.get_event();
+            }
+            else if (m_size != count) {
+               T* newData = nullptr;
+               hipMallocAsync(&newData, count * sizeof(T), m_control->m_last_resource.get_stream());
+               const size_t min = m_size < count ? m_size : count;
+               hipMemcpyAsync(newData, m_data, min * sizeof(T), hipMemcpyDeviceToDevice, m_control->m_last_resource);
+               m_control->m_last_event = m_control->m_last_resource.get_event();
+               hipFreeAsync(m_data, m_control->m_last_resource);
+               m_data = newData;
+            }
+
+            m_size = count;
+         }
+
+         void free() {
+            if (m_control) {
+               hipFreeAsync(m_data, m_control->m_last_resource.get_stream());
+               delete m_control;
+               m_control = nullptr;
+               m_data = nullptr;
+               m_size = 0;
+            }
+
+         T& operator[](size_t i) const {
+            return m_data[i];
+         }
+
+      private:
+         struct ControlBlock {
+            ControlBlock(size_t count,
+                         T* data,
+                         camp::resources::Resource resource) :
+               m_count{count},
+               m_data{data},
+               m_resource{resource},
+               m_event{resource.get_event()}
+            {}
+
+            size_t m_count = 0;
+            T* m_data = nullptr;
+            camp::resources::Resource m_last_resource;
+            camp::resources::Event m_last_event;
+         };
+
+         size_t m_size = 0;
+         T* m_data = nullptr;
+         ControlBlock* m_control = nullptr;
+   };
+
+}  // namespace care
+
+#endif  // CARE_PINNED_ARRAY_HPP
diff --git a/src/chai/PinnedPlugin.cpp b/src/chai/PinnedPlugin.cpp
new file mode 100644
index 00000000..02453c7e
--- /dev/null
+++ b/src/chai/PinnedPlugin.cpp
@@ -0,0 +1,74 @@
+//////////////////////////////////////////////////////////////////////////////
+// Copyright (c) 2016-24, Lawrence Livermore National Security, LLC and CHAI
+// project contributors. See the CHAI LICENSE file for details.
+//
+// SPDX-License-Identifier: BSD-3-Clause
+//////////////////////////////////////////////////////////////////////////////
+#include "chai/config.hpp"
+#include "chai/PinnedPlugin.hpp"
+
+namespace chai {
+
+RajaExecutionSpacePlugin::RajaExecutionSpacePlugin()
+{
+}
+
+void
+RajaExecutionSpacePlugin::preCapture(const RAJA::util::PluginContext& p)
+{
+  if (!m_arraymanager) {
+    m_arraymanager = chai::ArrayManager::getInstance();
+  }
+
+  switch (p.platform) {
+    case RAJA::Platform::host:
+      m_arraymanager->setExecutionSpace(chai::CPU); break;
+#if defined(CHAI_ENABLE_CUDA)
+    case RAJA::Platform::cuda:
+      m_arraymanager->setExecutionSpace(chai::GPU); break;
+#endif
+#if defined(CHAI_ENABLE_HIP)
+    case RAJA::Platform::hip:
+      m_arraymanager->setExecutionSpace(chai::GPU); break;
+#endif
+    default:
+      m_arraymanager->setExecutionSpace(chai::NONE);
+  }
+}
+
+void
+RajaExecutionSpacePlugin::postCapture(const RAJA::util::PluginContext&)
+{
+  m_arraymanager->setExecutionSpace(chai::NONE);
+}
+
+}
+RAJA_INSTANTIATE_REGISTRY(RAJA::util::PluginRegistry);
+
+// this is needed to link a dynamic lib as RAJA does not provide an exported definition of this symbol.
+#if defined(_WIN32) && !defined(CHAISTATICLIB)
+#ifdef CHAISHAREDDLL_EXPORTS
+namespace RAJA
+{
+namespace util
+{
+
+PluginStrategy::PluginStrategy() = default;
+
+}  // namespace util
+}  // namespace RAJA
+#endif
+#endif
+
+// Register plugin with RAJA
+RAJA::util::PluginRegistry::add<chai::RajaExecutionSpacePlugin> P(
+     "RajaExecutionSpacePlugin",
+     "Plugin to set CHAI execution space based on RAJA execution platform");
+
+
+namespace chai {
+
+  void linkRajaPlugin() {}
+
+}
+
diff --git a/src/chai/PinnedPlugin.hpp b/src/chai/PinnedPlugin.hpp
new file mode 100644
index 00000000..19152af9
--- /dev/null
+++ b/src/chai/PinnedPlugin.hpp
@@ -0,0 +1,29 @@
+//////////////////////////////////////////////////////////////////////////////
+// Copyright (c) 2016-24, Lawrence Livermore National Security, LLC and CHAI
+// project contributors. See the CHAI LICENSE file for details.
+//
+// SPDX-License-Identifier: BSD-3-Clause
+//////////////////////////////////////////////////////////////////////////////
+#if !defined(CHAI_PINNED_PLUGIN_HPP)
+#define CHAI_PINNED_PLUGIN_HPP
+
+#include "RAJA/util/PluginStrategy.hpp"
+
+namespace chai {
+
+class PinnedPlugin : public RAJA::util::PluginStrategy
+{
+  public:
+    PinnedPlugin();
+
+    void preCapture(const RAJA::util::PluginContext& p) override;
+
+    void postCapture(const RAJA::util::PluginContext& p) override;
+
+    void postLaunch(const RAJA::util::PluginContext& p) override;
+
+  private:
+    // Some list of callbacks
+};
+
+#endif  // CHAI_PINNED_PLUGIN_HPP
diff --git a/src/chai/containers/Allocator.hpp b/src/chai/containers/Allocator.hpp
new file mode 100644
index 00000000..ce26b4d1
--- /dev/null
+++ b/src/chai/containers/Allocator.hpp
@@ -0,0 +1,44 @@
+#ifndef CHAI_ALLOCATOR_HPP
+#define CHAI_ALLOCATOR_HPP
+
+#include "chai/MemoryType.hpp"
+
+namespace chai {
+   class Allocator {
+      public:
+         static void setAllocator(MemoryType type, int id) {
+            const umpire::Allocator& allocator =
+               umpire::ResourceManager::getInstance().getAllocator(id);
+
+            switch (type) {
+               case MemoryType::Host:
+                  getHostAllocator() = allocator;
+                  break;
+            }
+         }
+
+         void* allocate(MemoryType type, size_t size) {
+            switch (type) {
+               case MemoryType::Host:
+                  return getHostAllocator().allocate(size);
+            }
+         }
+
+         void deallocate(MemoryType type, void* data) {
+            switch (type) {
+               case MemoryType::Host:
+                  getHostAllocator().deallocate(data);
+            }
+         }
+
+      private:
+         static umpire::Allocator& getHostAllocator() {
+            static umpire::Allocator s_hostAllocator =
+               umpire::ResourceManager::getInstance().getAllocator("HOST");
+
+            return s_hostAllocator;
+         }
+   };  // namespace Allocator
+}  // namespace chai
+
+#endif  // CHAI_ALLOCATOR_HPP
diff --git a/src/chai/containers/CopyHidingManager.hpp b/src/chai/containers/CopyHidingManager.hpp
new file mode 100644
index 00000000..45db8c8a
--- /dev/null
+++ b/src/chai/containers/CopyHidingManager.hpp
@@ -0,0 +1,68 @@
+#ifndef CHAI_COPY_HIDING_MANAGER_HPP
+#define CHAI_COPY_HIDING_MANAGER_HPP
+
+namespace chai {
+   template <class T, class Allocator>
+   class CopyHidingManager {
+      public:
+         ~CopyHidingManager()
+         {
+            m_allocator.deallocate(chai::GPU, m_gpu_data);
+            m_allocator.deallocate(chai::CPU, m_cpu_data);
+         }
+
+         T* allocate(size_t size)
+         {
+            m_gpu_data = m_allocator.allocate(chai::GPU, size);
+            m_cpu_data = m_allocator.allocate(chai::CPU, size);
+         }
+
+         T* reallocate(size_t size)
+         {
+            m_gpu_data = m_allocator.reallocate(chai::GPU, size);
+            m_cpu_data = m_allocator.reallocate(chai::CPU, size);
+         }
+
+         T* data(ExecutionSpace executionSpace,
+                 bool update = true,
+                 bool touch = true)
+         {
+            if (update && executionSpace != m_space)
+            {
+
+            }
+
+            if (touch)
+            {
+               m_space = executionSpace;
+            }
+            else
+            {
+               m_space = NONE;
+            }
+
+            if (executionSpace == CPU)
+            {
+               return m_cpu_data;
+            }
+            else if (executionSpace == GPU)
+            {
+               return m_gpu_data;
+            }
+            else
+            {
+               return nullptr;
+            }
+         }
+
+      private:
+         T* m_cpu_data = nullptr;
+         T* m_gpu_data = nullptr;
+         size_t m_size = 0;
+         chai::ExecutionSpace m_space = NONE;
+         Allocator m_allocator;
+
+   };
+}  // namespace chai
+
+#endif  // CHAI_COPY_HIDING_MANAGER_HPP
diff --git a/src/chai/containers/DeviceMemoryResource.hpp b/src/chai/containers/DeviceMemoryResource.hpp
new file mode 100644
index 00000000..648a6fd8
--- /dev/null
+++ b/src/chai/containers/DeviceMemoryResource.hpp
@@ -0,0 +1,50 @@
+#ifndef CHAI_DEVICE_MEMORY_RESOURCE_HPP
+#define CHAI_DEVICE_MEMORY_RESOURCE_HPP
+
+#include "chai/MemoryType.hpp"
+
+namespace chai {
+   template <class AllocatorType>
+   class DeviceMemoryResource {
+      public:
+         ~DeviceMemoryResource() {
+            deallocate();
+         }
+
+         void* allocate(size_t size) {
+            m_data = m_allocator.allocate(MemoryType::Device, size);
+            m_size = size;
+            return m_data;
+         }
+
+         void* reallocate(size_t size) {
+            void* newData = m_allocator.allocate(MemoryType::Device, size);
+            memcpy(newData, m_data, size < m_size ? size : m_size);
+            m_allocator.deallocate(MemoryType::Device, m_data);
+            m_data = newData;
+            m_size = size;
+            return m_data;
+         }
+
+         void deallocate() {
+            m_allocator.deallocate(MemoryType::Device, m_data);
+            m_data = nullptr;
+            m_size = 0;
+         }
+
+         void* data(bool touch) {
+            return data(MemoryResourcePlugin::getExecutionSpace(), touch);
+         }
+
+         void* data(ExecutionSpace /* executionSpace */, bool /* touch */) {
+            return m_data;
+         }
+
+      private:
+         void* m_data = nullptr;
+         size_t m_size = 0;
+         AllocatorType m_allocator;
+   };  // class DeviceMemoryResource
+}  // namespace chai
+
+#endif  // CHAI_DEVICE_MEMORY_RESOURCE_HPP
diff --git a/src/chai/containers/HostMemoryResource.hpp b/src/chai/containers/HostMemoryResource.hpp
new file mode 100644
index 00000000..5950364b
--- /dev/null
+++ b/src/chai/containers/HostMemoryResource.hpp
@@ -0,0 +1,51 @@
+#ifndef CHAI_HOST_MEMORY_RESOURCE_HPP
+#define CHAI_HOST_MEMORY_RESOURCE_HPP
+
+#include "chai/MemoryType.hpp"
+
+namespace chai {
+   template <class AllocatorType>
+   class HostMemoryResource {
+      public:
+         ~HostMemoryResource() {
+            deallocate();
+         }
+
+         void* allocate(size_t size) {
+            m_data = m_allocator.allocate(MemoryType::Host, size);
+            m_size = size;
+            return m_data;
+         }
+
+         void* reallocate(size_t size) {
+            void* newData = m_allocator.allocate(MemoryType::Host, size);
+            memcpy(newData, m_data, size < m_size ? size : m_size);
+            m_allocator.deallocate(MemoryType::Host, m_data);
+            m_data = newData;
+            m_size = size;
+            return m_data;
+         }
+
+         void deallocate() {
+            m_allocator.deallocate(MemoryType::Host, m_data);
+            m_data = nullptr;
+            m_size = 0;
+         }
+
+         void* data(bool touch) {
+            return data(MemoryResourcePlugin::getExecutionSpace(), touch);
+         }
+
+         void* data(ExecutionSpace /* executionSpace */, bool /* touch */) {
+            return m_data;
+         }
+
+      private:
+         void* m_data = nullptr;
+         size_t m_size = 0;
+         AllocatorType m_allocator;
+   };  // class HostMemoryResource
+
+}  // namespace chai
+
+#endif  // CHAI_HOST_MEMORY_RESOURCE_HPP
diff --git a/src/chai/containers/MemoryType.hpp b/src/chai/containers/MemoryType.hpp
new file mode 100644
index 00000000..75a0d603
--- /dev/null
+++ b/src/chai/containers/MemoryType.hpp
@@ -0,0 +1,8 @@
+#ifndef CHAI_MEMORY_TYPE_HPP
+#define CHAI_MEMORY_TYPE_HPP
+
+namespace chai {
+   enum class MemoryType { Host, Device };
+}  // namespace chai
+
+#endif  // CHAI_MEMORY_TYPE_HPP
diff --git a/src/chai/containers/PArray.hpp b/src/chai/containers/PArray.hpp
new file mode 100644
index 00000000..76f71e7a
--- /dev/null
+++ b/src/chai/containers/PArray.hpp
@@ -0,0 +1,109 @@
+#ifndef CHAI_PARRAY_HPP
+#define CHAI_PARRAY_HPP
+
+#include "chai/config.hpp"
+
+namespace chai {
+   template <class T, class Manager>
+   class PArray
+   {
+      public:
+         constexpr PArray() = default;
+
+         PArray(size_t size)
+         {
+            if (size > 0)
+            {
+               allocate(size);
+            }
+         }
+
+         CHAI_HOST_DEVICE PArray(const PArray& 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();
+            }
+#endif
+         }
+
+         CHAI_HOST_DEVICE PArray& operator=(const PArray& other) = default;
+
+         void reallocate(size_t size)
+         {
+            if (size == 0)
+            {
+               deallocate();
+            }
+            else if (m_size == 0)
+            {
+               allocate();
+            }
+            else
+            {
+               m_data = m_manager->reallocate(size);
+               m_size = size;
+            }
+         }
+
+         void deallocate()
+         {
+            if (m_manager) {
+               delete m_manager;
+               m_manager = nullptr;
+               m_data = nullptr;
+               m_size = 0;
+            }
+         }
+
+         CHAI_HOST_DEVICE size_t size() const
+         {
+            return m_size;
+         }
+
+         CHAI_HOST_DEVICE T& operator[](size_t i) const
+         {
+            return m_data[i];
+         }
+
+         CHAI_HOST_DEVICE T* data() const
+         {
+#if !defined(CHAI_DEVICE_COMPILE)
+            if (m_manager)
+            {
+               m_data = m_manager->data();
+            }
+#endif
+            return m_data;
+         }
+
+         CHAI_HOST_DEVICE const T* cdata() const
+         {
+#if !defined(CHAI_DEVICE_COMPILE)
+            if (m_manager)
+            {
+               m_data = m_manager->cdata();
+            }
+#endif
+            return m_data;
+         }
+
+      private:
+         T* m_data = nullptr;
+         size_t m_size = 0;
+         Manager* m_manager = nullptr;
+
+         void allocate(size_t size)
+         {
+            m_manager = new Manager();
+            m_data = m_manager->allocate(size);
+            m_size = size;
+         }
+   };  // class PArray
+}  // namespace chai
+
+#endif  // CHAI_PARRAY_HPP