ManagedSharedPtr

src/chai/ArrayManager.hpp

@@ -7,122 +7,10 @@
 #ifndef CHAI_ArrayManager_HPP
 #define CHAI_ArrayManager_HPP
 
-#include "chai/config.hpp"
-#include "chai/ChaiMacros.hpp"
-#include "chai/ExecutionSpaces.hpp"
-#include "chai/PointerRecord.hpp"
-#include "chai/Types.hpp"
-
-#if defined(CHAI_ENABLE_RAJA_PLUGIN)
-#include "chai/pluginLinker.hpp"
-#endif
-
-#include <unordered_map>
-
-#include "umpire/Allocator.hpp"
-#include "umpire/util/MemoryMap.hpp"
-
-#if defined(CHAI_ENABLE_CUDA)
-#include <cuda_runtime_api.h>
-#endif
-#if defined(CHAI_ENABLE_HIP)
-#include "hip/hip_runtime_api.h"
-#endif
+#include "chai/ChaiManager.hpp"
 
 namespace chai
 {
-// CHAI_GPU_ERROR_CHECK macro
-#if defined(CHAI_ENABLE_CUDA) || defined(CHAI_ENABLE_HIP)
-
-#ifdef CHAI_ENABLE_GPU_ERROR_CHECKING
-
-#ifdef CHAI_ENABLE_CUDA
-inline void gpuErrorCheck(cudaError_t code, const char *file, int line, bool abort=true)
-{
-   if (code != cudaSuccess) {
-      fprintf(stderr, "[CHAI] GPU Error: %s %s %d\n", cudaGetErrorString(code), file, line);
-      if (abort) {
-         exit(code);
-      }
-   }
-}
-#elif defined(CHAI_ENABLE_HIP)
-inline void gpuErrorCheck(hipError_t code, const char *file, int line, bool abort=true)
-{
-   if (code != hipSuccess) {
-      fprintf(stderr, "[CHAI] GPU Error: %s %s %d\n", hipGetErrorString(code), file, line);
-      if (abort) {
-         exit(code);
-      }
-   }
-}
-#endif
-
-
-#define CHAI_GPU_ERROR_CHECK(code) { gpuErrorCheck((code), __FILE__, __LINE__); }
-#else // CHAI_ENABLE_GPU_ERROR_CHECKING
-#define CHAI_GPU_ERROR_CHECK(code) code
-#endif // CHAI_ENABLE_GPU_ERROR_CHECKING
-
-#endif
-
-// wrapper for hip/cuda synchronize
-inline void synchronize() {
-#if defined (CHAI_ENABLE_HIP) &&!defined(__HIP_DEVICE_COMPILE__)
-   CHAI_GPU_ERROR_CHECK(hipDeviceSynchronize());
-#elif defined (CHAI_ENABLE_CUDA) &&!defined(__CUDA_ARCH__)
-   CHAI_GPU_ERROR_CHECK(cudaDeviceSynchronize());
-#endif
-}
-
-#if defined(CHAI_GPUCC) || defined(CHAI_ENABLE_GPU_SIMULATION_MODE)
-
-// wrapper for hip/cuda free
-CHAI_HOST inline void gpuFree(void* buffer) {
-#if defined(CHAI_ENABLE_GPU_SIMULATION_MODE)
-   free(buffer);
-#elif defined (CHAI_ENABLE_HIP)
-   CHAI_GPU_ERROR_CHECK(hipFree(buffer));
-#elif defined (CHAI_ENABLE_CUDA)
-   CHAI_GPU_ERROR_CHECK(cudaFree(buffer));
-#endif
-}
-
-// wrapper for hip/cuda malloc
-CHAI_HOST inline void gpuMalloc(void** devPtr, size_t size) {
-#if defined(CHAI_ENABLE_GPU_SIMULATION_MODE)
-   *devPtr = (void*)malloc(size);
-#elif defined (CHAI_ENABLE_HIP)
-   CHAI_GPU_ERROR_CHECK(hipMalloc(devPtr, size));
-#elif defined (CHAI_ENABLE_CUDA)
-   CHAI_GPU_ERROR_CHECK(cudaMalloc(devPtr, size));
-#endif
-}
-
-// wrapper for hip/cuda managed malloc
-CHAI_HOST inline void gpuMallocManaged(void** devPtr, size_t size) {
-#if defined(CHAI_ENABLE_GPU_SIMULATION_MODE)
-   *devPtr = (void*)malloc(size);
-#elif defined (CHAI_ENABLE_HIP)
-   CHAI_GPU_ERROR_CHECK(hipMallocManaged(devPtr, size));
-#elif defined (CHAI_ENABLE_CUDA)
-   CHAI_GPU_ERROR_CHECK(cudaMallocManaged(devPtr, size));
-#endif
-}
-
-// wrapper for hip/cuda mem copy
-CHAI_HOST inline void  gpuMemcpy(void* dst, const void* src, size_t count, gpuMemcpyKind kind) {
-#if defined(CHAI_ENABLE_GPU_SIMULATION_MODE)
-   memcpy(dst, src, count);
-#elif defined (CHAI_ENABLE_HIP)
-   CHAI_GPU_ERROR_CHECK(hipMemcpy(dst, src, count, kind));
-#elif defined (CHAI_ENABLE_CUDA)
-   CHAI_GPU_ERROR_CHECK(cudaMemcpy(dst, src, count, kind));
-#endif
-}
-
-#endif //#if defined(CHAI_GPUCC)
-
 /*!
  * \brief Singleton that manages caching and movement of ManagedArray objects.
  *

src/chai/CMakeLists.txt

@@ -15,6 +15,7 @@ set (chai_headers
   ArrayManager.hpp
   ArrayManager.inl
   ChaiMacros.hpp
+  ChaiManager.hpp
   ExecutionSpaces.hpp
   ManagedArray.hpp
   ManagedArray.inl
@@ -29,6 +30,7 @@ if(CHAI_DISABLE_RM)
 endif ()
 
 set (chai_sources
+  SharedPtrManager.cpp
   ArrayManager.cpp)
 
 set (chai_depends

src/chai/ChaiMacros.hpp

@@ -31,6 +31,11 @@
 #define gpuMemcpyDeviceToDevice cudaMemcpyDeviceToDevice
 #define gpuMemcpyDefault cudaMemcpyDefault
 
+#define gpuSuccess cudaSuccess
+#define gpuError_t cudaError_t
+#define gpuGetErrorString cudaGetErrorString
+#define gpuDeviceSynchronize cudaDeviceSynchronize
+
 // NOTE: Cannot have if defined(__HIPCC__) in the condition below, since __HIPCC__ comes from the included header hip_runtime below.
 #elif defined(CHAI_ENABLE_HIP)
 
@@ -48,6 +53,11 @@
 #define gpuMemcpyDeviceToDevice hipMemcpyDeviceToDevice
 #define gpuMemcpyDefault hipMemcpyDefault
 
+#define gpuSuccess hipSuccess
+#define gpuError_t hipError_t
+#define gpuGetErrorString hipGetErrorString
+#define gpuDeviceSynchronize hipDeviceSynchronize
+
 #else
 
 #define CHAI_HOST

src/chai/ChaiManager.hpp

@@ -0,0 +1,28 @@
+//////////////////////////////////////////////////////////////////////////////
+// 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
+//////////////////////////////////////////////////////////////////////////////
+#ifndef CHAI_ChaiManager_HPP
+#define CHAI_ChaiManager_HPP
+
+#include "chai/ChaiMacros.hpp"
+#include "chai/ExecutionSpaces.hpp"
+#include "chai/Types.hpp"
+
+#include "chai/PointerRecord.hpp"
+
+#if defined(CHAI_ENABLE_RAJA_PLUGIN)
+#include "chai/pluginLinker.hpp"
+#endif
+
+#include <unordered_map>
+
+#include "umpire/Allocator.hpp"
+#include "umpire/util/MemoryMap.hpp"
+
+
+#include "chai/util/DeviceHelpers.hpp"
+
+#endif // CHAI_ChaiManager_HPP

src/chai/ManagedArray.hpp

@@ -373,14 +373,14 @@ class ManagedArray : public CHAICopyable
   // shenanigan reasons need to be defined here.
 #if !defined(CHAI_DISABLE_RM)
   // if T is a CHAICopyable, then it is important to initialize all the
-  // ManagedArrays to nullptr at allocation, since it is extremely easy to
+  // elements with default constructors, since it is extremely easy to
   // trigger a moveInnerImpl, which expects inner values to be initialized.
   template <bool B = std::is_base_of<CHAICopyable, T>::value,
             typename std::enable_if<B, int>::type = 0>
   CHAI_HOST bool initInner(size_t start = 0)
   {
     for (size_t i = start; i < m_size/sizeof(T); ++i) {
-      m_active_base_pointer[i] = nullptr;
+      new (&m_active_base_pointer[i]) T();
     }
     return true;
   }
@@ -392,6 +392,26 @@ class ManagedArray : public CHAICopyable
   {
     return false;
   }
+
+  // if T is a CHAICopyable, then it is important to free all the
+  // CHAICopyable containers, which expect inner values to be initialized.
+  template <bool B = std::is_base_of<CHAICopyable, T>::value,
+            typename std::enable_if<B, int>::type = 0>
+  CHAI_HOST bool freeInner(size_t start = 0)
+  {
+    for (size_t i = start; i < m_size/sizeof(T); ++i) {
+      m_active_base_pointer[i].~T();
+    }
+    return true;
+  }
+
+  // Do not deep initialize if T is not a CHAICopyable.
+  template <bool B = std::is_base_of<CHAICopyable, T>::value,
+            typename std::enable_if<!B, int>::type = 0>
+  CHAI_HOST bool freeInner(size_t = 0)
+  {
+    return false;
+  }
 #endif
 protected:
   /*!

src/chai/ManagedArray.inl

@@ -249,12 +249,14 @@ CHAI_HOST void ManagedArray<T>::reallocate(size_t elems)
       // trigger a moveInnerImpl, which expects inner values to be initialized.
       if (initInner(old_size/sizeof(T))) {
         // if we are active on the  GPU, we need to send any newly initialized inner members to the device
+#if defined(CHAI_ENABLE_CUDA) || defined(CHAI_ENABLE_HIP)
         if (m_pointer_record->m_last_space == GPU && old_size < m_size) {
           umpire::ResourceManager & umpire_rm = umpire::ResourceManager::getInstance();
           void *src = (void *)(((char *)(m_pointer_record->m_pointers[CPU])) + old_size);
           void *dst = (void *)(((char *)(m_pointer_record->m_pointers[GPU])) + old_size);
           umpire_rm.copy(dst,src,m_size-old_size);
         }
+#endif
       }
 
       CHAI_LOG(Debug, "m_active_ptr reallocated at address: " << m_active_pointer);
@@ -276,6 +278,8 @@ CHAI_HOST void ManagedArray<T>::free(ExecutionSpace space)
     if (m_pointer_record == &ArrayManager::s_null_record) {
        m_pointer_record = m_resource_manager->makeManaged((void *)m_active_base_pointer,m_size,space,true);
     }
+    freeInner();
+
     m_resource_manager->free(m_pointer_record, space);
     m_active_pointer = nullptr;
     m_active_base_pointer = nullptr;
@@ -302,6 +306,9 @@ CHAI_HOST void ManagedArray<T>::reset()
 template<typename T>
 CHAI_INLINE
 CHAI_HOST_DEVICE size_t ManagedArray<T>::size() const {
+#if !defined(CHAI_DEVICE_COMPILE)
+  if (!m_is_slice) m_size = m_pointer_record->m_size;
+#endif
   return m_size/sizeof(T);
 }