From f21c17f445ee0b92fca763699f4f929d54f7d2f6 Mon Sep 17 00:00:00 2001 From: Max Yang Date: Tue, 11 Feb 2020 12:51:59 -0800 Subject: [PATCH 01/70] Don't overwrite ownership info in registerPointer if a pointer has been set --- src/chai/ArrayManager.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/chai/ArrayManager.cpp b/src/chai/ArrayManager.cpp index bb91d71b..87a233e6 100644 --- a/src/chai/ArrayManager.cpp +++ b/src/chai/ArrayManager.cpp @@ -58,7 +58,7 @@ void ArrayManager::registerPointer( //record->m_last_space = space; for (int i = 0; i < NUM_EXECUTION_SPACES; i++) { - record->m_owned[i] = true; + if (!record->m_pointers[i]) record->m_owned[i] = true; } record->m_owned[space] = owned; } From 19759dda07ccea63bd6f470ea1faca621550d82a Mon Sep 17 00:00:00 2001 From: Max Yang Date: Tue, 11 Feb 2020 12:53:37 -0800 Subject: [PATCH 02/70] Set PointerRecord allocators in makeManagedArray --- src/chai/ManagedArray.hpp | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/src/chai/ManagedArray.hpp b/src/chai/ManagedArray.hpp index 37d16e91..ab30a868 100644 --- a/src/chai/ManagedArray.hpp +++ b/src/chai/ManagedArray.hpp @@ -358,6 +358,11 @@ ManagedArray makeManagedArray(T* data, PointerRecord* record = manager->makeManaged(data, sizeof(T) * elems, space, owned); + for (int space = CPU; space < NUM_EXECUTION_SPACES; space++) { + record->m_allocators[space] = + manager->getAllocatorId(ExecutionSpace(space)); + } + ManagedArray array = ManagedArray(record, space); if (!std::is_const::value) { From c6714922c8dd6f10c4ebb52867a6a4ac015aa55a Mon Sep 17 00:00:00 2001 From: Max Yang Date: Wed, 12 Feb 2020 11:02:45 -0800 Subject: [PATCH 03/70] Don't try to move nested array to GPU if already on GPU --- src/chai/ManagedArray.inl | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/chai/ManagedArray.inl b/src/chai/ManagedArray.inl index 87cb6d36..d3b809d5 100644 --- a/src/chai/ManagedArray.inl +++ b/src/chai/ManagedArray.inl @@ -343,7 +343,7 @@ void ManagedArray::move(ExecutionSpace space) /* When moving from GPU to CPU we need to move the inner arrays after the outer array. */ #if defined(CHAI_ENABLE_CUDA) - if (prev_space == GPU) { + if (space != GPU && prev_space == GPU) { moveInnerImpl(space); } #endif From 593a3c9f3175251f79bfc360fb7febd91d9c5b90 Mon Sep 17 00:00:00 2001 From: Max Yang Date: Wed, 12 Feb 2020 11:28:39 -0800 Subject: [PATCH 04/70] Add tests for makeManagedArray, nested array double move --- tests/integration/managed_array_tests.cpp | 84 +++++++++++++++++++++++ 1 file changed, 84 insertions(+) diff --git a/tests/integration/managed_array_tests.cpp b/tests/integration/managed_array_tests.cpp index 506b7c70..95d40bea 100644 --- a/tests/integration/managed_array_tests.cpp +++ b/tests/integration/managed_array_tests.cpp @@ -734,6 +734,24 @@ TEST(ManagedArray, ExternalConstructorOwned) array.free(); } +#if defined(CHAI_ENABLE_CUDA) || defined(CHAI_ENABLE_HIP) +GPU_TEST(ManagedArray, ExternalUnownedMoveToGPU) +{ + float data[20]; + for (int i = 0; i < 20; i++) { + data[i] = 0.; + } + + chai::ManagedArray array = + chai::makeManagedArray(data, 20, chai::CPU, false); + + forall(gpu(), 0, 20, [=] __device__ (int i) { array[i] = 1.0f * i; }); + + forall(sequential(), 0, 20, [=] (int i) { ASSERT_EQ(array[i], 1.0f * i); }); + + array.free(); +} +#endif #endif TEST(ManagedArray, Reset) @@ -1281,6 +1299,72 @@ GPU_TEST(ManagedArray, MoveInnerToDevice2) outerArray.free(); } +GPU_TEST(ManagedArray, MoveInnerToDeviceAgain) +{ + const int N = 5; + + /* Create the outer array. */ + chai::ManagedArray> outerArray(N); + + /* Loop over the outer array and populate it with arrays on the CPU. */ + forall(sequential(), 0, N, + [=](int i) + { + chai::ManagedArray temp(N); + + forall(sequential(), 0, N, + [=](int j) + { + temp[j] = N * i + j; + } + ); + + outerArray[i] = temp; + } + ); + + /* Capture the outer array on the GPU and update the values of the inner + * arrays. */ + forall(gpu(), 0, N, + [=] __device__(int i) + { + for( int j = 0; j < N; ++j) + { + outerArray[i][j] *= 2; + } + } + ); + + /* Capture the outer array on the GPU and update the values of the inner + * arrays. This time, the array should already be resident on the GPU. */ + forall(gpu(), 0, N, + [=] __device__(int i) + { + for( int j = 0; j < N; ++j) + { + outerArray[i][j] *= 2; + } + } + ); + + /* Capture the outer array on the CPU and check the values of the inner + * arrays. */ + forall(sequential(), 0, N, + [=](int i) + { + for (int j = 0; j < N; ++j) + { + ASSERT_EQ(outerArray[i][j], 4 * (N * i + j)); + } + } + ); + + for (int i = 0; i < N; ++i) { + outerArray[i].free(); + } + + outerArray.free(); +} #endif // CHAI_DISABLE_RM #endif // defined(CHAI_ENABLE_CUDA) || defined(CHAI_ENABLE_HIP) From 39785b26f6c843e63b55e8dcc156277878b698a9 Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Wed, 1 Apr 2020 08:03:12 -0700 Subject: [PATCH 05/70] BLT configuration fixes. Use proper library names, track Umpires possible dependency on mpi with ENABLE_MPI. --- CMakeLists.txt | 1 + cmake/thirdparty/SetupChaiThirdparty.cmake | 13 +++++++++++-- 2 files changed, 12 insertions(+), 2 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index a145df58..5e6b47e9 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -12,6 +12,7 @@ project(Chai LANGUAGES CXX VERSION 2.0.0) set(ENABLE_CUDA Off CACHE BOOL "Enable CUDA") set(ENABLE_HIP Off CACHE BOOL "Enable HIP") set(ENABLE_OPENMP On CACHE BOOL "Enable OpenMP") +set(ENABLE_MPI Off CACHE BOOL "Enable MPI (for umpire replay only)") set(ENABLE_BENCHMARKS On CACHE BOOL "Enable benchmarks") option(ENABLE_IMPLICIT_CONVERSIONS "Enable implicit conversions to-from raw pointers" On) option(DISABLE_RM "Make ManagedArray a thin wrapper" Off) diff --git a/cmake/thirdparty/SetupChaiThirdparty.cmake b/cmake/thirdparty/SetupChaiThirdparty.cmake index 21a63f9a..ba3246a4 100644 --- a/cmake/thirdparty/SetupChaiThirdparty.cmake +++ b/cmake/thirdparty/SetupChaiThirdparty.cmake @@ -7,11 +7,16 @@ if (NOT TARGET umpire) if (DEFINED umpire_DIR) find_package(umpire REQUIRED) - + if (ENABLE_MPI) + set(UMPIRE_DEPENDS mpi) + else() + set(UMPIRE_DEPENDS) + endif() blt_register_library( NAME umpire INCLUDES ${UMPIRE_INCLUDE_DIRS} - LIBRARIES umpire) + LIBRARIES ${UMPIRE_LIBNAME} + DEPENDS_ON ${UMPIRE_DEPENDS}) else () set(OLD_ENABLE_FORTRAN ${ENABLE_FORTRAN}) set(ENABLE_FORTRAN Off CACHE BOOL "Enable Fortran in Umpire") @@ -25,6 +30,10 @@ if (ENABLE_RAJA_PLUGIN) if (DEFINED RAJA_DIR) message(STATUS "CHAI: using external RAJA via find_package") find_package(RAJA REQUIRED) + blt_register_library( + NAME RAJA + INCLUDES ${RAJA_INCLUDE_DIRS} + LIBRARIES ${RAJA_LIBRARY}) else() message(STATUS "CHAI: using builtin RAJA submodule") add_subdirectory(${PROJECT_SOURCE_DIR}/src/tpl/raja) From 3bd43f748d8680ef129e92457f1a99a7ca66c171 Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Thu, 2 Apr 2020 10:50:46 -0700 Subject: [PATCH 06/70] switch back to using cmake targets --- cmake/thirdparty/SetupChaiThirdparty.cmake | 12 +++++++++--- 1 file changed, 9 insertions(+), 3 deletions(-) diff --git a/cmake/thirdparty/SetupChaiThirdparty.cmake b/cmake/thirdparty/SetupChaiThirdparty.cmake index ba3246a4..492a62d9 100644 --- a/cmake/thirdparty/SetupChaiThirdparty.cmake +++ b/cmake/thirdparty/SetupChaiThirdparty.cmake @@ -6,6 +6,9 @@ ############################################################################## if (NOT TARGET umpire) if (DEFINED umpire_DIR) + # this allows umpire_DIR to be the install prefix if we are relying on umpire's + # installed umpire-config.cmake + list(APPEND CMAKE_PREFIX_PATH ${umpire_DIR}) find_package(umpire REQUIRED) if (ENABLE_MPI) set(UMPIRE_DEPENDS mpi) @@ -15,7 +18,7 @@ if (NOT TARGET umpire) blt_register_library( NAME umpire INCLUDES ${UMPIRE_INCLUDE_DIRS} - LIBRARIES ${UMPIRE_LIBNAME} + LIBRARIES umpire DEPENDS_ON ${UMPIRE_DEPENDS}) else () set(OLD_ENABLE_FORTRAN ${ENABLE_FORTRAN}) @@ -28,12 +31,15 @@ endif() if (ENABLE_RAJA_PLUGIN) if (NOT TARGET RAJA) if (DEFINED RAJA_DIR) - message(STATUS "CHAI: using external RAJA via find_package") + # this allows RAJA_DIR to be the install prefix if we are relying on RAJA's + # installed RAJA-config.cmake + list(APPEND CMAKE_PREFIX_PATH ${RAJA_DIR}) + message(STATUS "CHAI: using external RAJA via find_package ${RAJA_DIR}") find_package(RAJA REQUIRED) blt_register_library( NAME RAJA INCLUDES ${RAJA_INCLUDE_DIRS} - LIBRARIES ${RAJA_LIBRARY}) + LIBRARIES RAJA) else() message(STATUS "CHAI: using builtin RAJA submodule") add_subdirectory(${PROJECT_SOURCE_DIR}/src/tpl/raja) From c837f20e48331f0fecaa548c43dbd958273cab38 Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Thu, 2 Apr 2020 11:31:54 -0700 Subject: [PATCH 07/70] adds ability to simulate GPU memory motion on a host-only build. --- src/chai/ArrayManager.cpp | 9 ++++++++- src/chai/ArrayManager.hpp | 8 ++++++++ src/chai/ArrayManager.inl | 4 ++++ src/chai/CMakeLists.txt | 1 + src/chai/ExecutionSpaces.hpp | 8 +++++++- src/chai/config.hpp.in | 1 + 6 files changed, 29 insertions(+), 2 deletions(-) diff --git a/src/chai/ArrayManager.cpp b/src/chai/ArrayManager.cpp index 87a233e6..71d0d1d6 100644 --- a/src/chai/ArrayManager.cpp +++ b/src/chai/ArrayManager.cpp @@ -34,10 +34,17 @@ ArrayManager::ArrayManager() : m_allocators[CPU] = new umpire::Allocator(m_resource_manager.getAllocator("HOST")); -#if defined(CHAI_ENABLE_CUDA) || defined(CHAI_ENABLE_HIP) + +#if defined(CHAI_ENABLE_CUDA) || defined(CHAI_ENABLE_HIP) || defined(CHAI_ENABLE_GPU_SIMULATION_MODE) +#if defined(CHAI_ENABLE_GPU_SIMULATION_MODE) + m_allocators[GPU] = + new umpire::Allocator(m_resource_manager.getAllocator("HOST")); +#else m_allocators[GPU] = new umpire::Allocator(m_resource_manager.getAllocator("DEVICE")); #endif +#endif + #if defined(CHAI_ENABLE_UM) m_allocators[UM] = new umpire::Allocator(m_resource_manager.getAllocator("UM")); diff --git a/src/chai/ArrayManager.hpp b/src/chai/ArrayManager.hpp index ddca8220..9f4d52a1 100644 --- a/src/chai/ArrayManager.hpp +++ b/src/chai/ArrayManager.hpp @@ -279,6 +279,14 @@ class ArrayManager */ void deregisterPointer(PointerRecord* record); + /*! + * \brief set the allocator for an execution space. + * + * \param space Execution space to set the default allocator for. + * \param allocator The allocator to use for this space. Will be copied into chai. + */ + void setAllocator(ExecutionSpace space, umpire::Allocator &allocator); + /*! * \brief Move data in PointerRecord to the corresponding ExecutionSpace. * diff --git a/src/chai/ArrayManager.inl b/src/chai/ArrayManager.inl index 90d4f280..6a56ac82 100644 --- a/src/chai/ArrayManager.inl +++ b/src/chai/ArrayManager.inl @@ -89,6 +89,10 @@ void ArrayManager::set(T* dst_ptr, size_t index, const T& val) } #endif +CHAI_INLINE +void ArrayManager::setAllocator(ExecutionSpace space, umpire::Allocator &allocator) { + *m_allocators[space] = allocator; +} } // end of namespace chai #endif // CHAI_ArrayManager_INL diff --git a/src/chai/CMakeLists.txt b/src/chai/CMakeLists.txt index c25de521..027abf9d 100644 --- a/src/chai/CMakeLists.txt +++ b/src/chai/CMakeLists.txt @@ -11,6 +11,7 @@ set(CHAI_ENABLE_IMPLICIT_CONVERSIONS ${ENABLE_IMPLICIT_CONVERSIONS}) set(CHAI_DISABLE_RM ${DISABLE_RM}) set(CHAI_ENABLE_UM ${ENABLE_UM}) set(CHAI_ENABLE_RAJA_PLUGIN ${ENABLE_RAJA_PLUGIN}) +set(CHAI_ENABLE_GPU_SIMULATION_MODE ${ENABLE_GPU_SIMULATION_MODE}) configure_file( ${PROJECT_SOURCE_DIR}/src/chai/config.hpp.in diff --git a/src/chai/ExecutionSpaces.hpp b/src/chai/ExecutionSpaces.hpp index 02f1c889..9fa32177 100644 --- a/src/chai/ExecutionSpaces.hpp +++ b/src/chai/ExecutionSpaces.hpp @@ -20,7 +20,7 @@ enum ExecutionSpace { NONE = 0, /*! Executing in CPU space */ CPU, -#if defined(CHAI_ENABLE_CUDA) || defined(CHAI_ENABLE_HIP) +#if defined(CHAI_ENABLE_CUDA) || defined(CHAI_ENABLE_HIP) || defined(CHAI_ENABLE_GPU_SIMULATION_MODE) /*! Execution in GPU space */ GPU, #endif @@ -30,6 +30,12 @@ enum ExecutionSpace { // NUM_EXECUTION_SPACES should always be last! /*! Used to count total number of spaces */ NUM_EXECUTION_SPACES +#if !defined(CHAI_ENABLE_CUDA) && !defined(CHAI_ENABLE_HIP) && !defined(CHAI_ENABLE_GPU_SIMULATION_MODE) + ,GPU +#endif +#if !defined(CHAI_ENABLE_UM) + ,UM +#endif }; } // end of namespace chai diff --git a/src/chai/config.hpp.in b/src/chai/config.hpp.in index 64f5ec66..ae420f5e 100644 --- a/src/chai/config.hpp.in +++ b/src/chai/config.hpp.in @@ -14,5 +14,6 @@ #cmakedefine CHAI_DISABLE_RM #cmakedefine CHAI_ENABLE_UM #cmakedefine CHAI_ENABLE_RAJA_PLUGIN +#cmakedefine CHAI_ENABLE_GPU_SIMULATION_MODE #endif // CHAI_config_HPP From 730183c37a9568b78d986e0d57e8ea4211a1b48a Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Thu, 2 Apr 2020 11:49:05 -0700 Subject: [PATCH 08/70] Update advanced configuration documentation. --- CMakeLists.txt | 1 + docs/sphinx/advanced_configuration.rst | 19 +++++++++++++++---- 2 files changed, 16 insertions(+), 4 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 5e6b47e9..6731b652 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -11,6 +11,7 @@ project(Chai LANGUAGES CXX VERSION 2.0.0) set(ENABLE_CUDA Off CACHE BOOL "Enable CUDA") set(ENABLE_HIP Off CACHE BOOL "Enable HIP") +set(ENABLE_GPU_SIMULATION_MODE Off CACHE BOOL "Enable GPU Simulation Mode") set(ENABLE_OPENMP On CACHE BOOL "Enable OpenMP") set(ENABLE_MPI Off CACHE BOOL "Enable MPI (for umpire replay only)") set(ENABLE_BENCHMARKS On CACHE BOOL "Enable benchmarks") diff --git a/docs/sphinx/advanced_configuration.rst b/docs/sphinx/advanced_configuration.rst index cda7d59e..cec6e697 100644 --- a/docs/sphinx/advanced_configuration.rst +++ b/docs/sphinx/advanced_configuration.rst @@ -15,19 +15,30 @@ Here is a summary of the configuration options, their default value, and meaning =========================== ======== =============================================================================== Variable Default Meaning =========================== ======== =============================================================================== - ENABLE_CUDA On Enable CUDA support + ENABLE_CUDA Off Enable CUDA support. + ENABLE_HIP Off Enable HIP support. + ENABLE_GPU_SIMULATION_MODE Off Simulates GPU execution. ENABLE_UM Off Enable support for CUDA Unified Memory. ENABLE_CNEM Off Enable cnmem for GPU allocations ENABLE_IMPLICIT_CONVERSIONS On Enable implicit conversions between ManagedArray and raw pointers DISABLE_RM Off Disable the ArrayManager and make ManagedArray a thin wrapper around a pointer. - ENABLE_TESTING On Build test executables - ENABLE_BENCHMARKS On Build benchmark programs + ENABLE_TESTING On Build test executables. + ENABLE_BENCHMARKS On Build benchmark programs. =========================== ======== =============================================================================== These arguments are explained in more detail below: * ENABLE_CUDA - This option enables support for GPUs. If CHAI is built without CUDA support, + This option enables support for GPUs using CUDA. If CHAI is built without CUDA, HIP, or + GPU_SIMULATION_MODE support, then only the ``CPU`` execution space is available for use. + +* ENABLE_HIP + This option enables support for GPUs using HIP. If CHAI is built without CUDA, HIP, or + GPU_SIMULATION_MODE support, then only the ``CPU`` execution space is available for use. + +* ENABLE_GPU_SIMULATION_MODE + This option simulates GPU support by enableing the GPU execution space, backed by a HOST + umpire allocator. If CHAI is built without CUDA, HIP, or GPU_SIMULATION_MODE support, then only the ``CPU`` execution space is available for use. * ENABLE_UM From e06089a2b2f4076d4c0047bf1c1f9c4dc7f5e4a8 Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Thu, 2 Apr 2020 15:21:53 -0700 Subject: [PATCH 09/70] add capability to evict all chai managed data from one space to another. --- src/chai/ArrayManager.cpp | 132 ++++++++++++++++++++++++++------------ src/chai/ArrayManager.hpp | 15 ++++- 2 files changed, 104 insertions(+), 43 deletions(-) diff --git a/src/chai/ArrayManager.cpp b/src/chai/ArrayManager.cpp index 71d0d1d6..f193a05d 100644 --- a/src/chai/ArrayManager.cpp +++ b/src/chai/ArrayManager.cpp @@ -189,52 +189,54 @@ void ArrayManager::allocate( CHAI_LOG(Debug, "Allocated array at: " << pointer_record->m_pointers[space]); } -void ArrayManager::free(PointerRecord* pointer_record) +void ArrayManager::free(PointerRecord* pointer_record, ExecutionSpace spaceToFree) { if (!pointer_record) return; for (int space = CPU; space < NUM_EXECUTION_SPACES; ++space) { - if (pointer_record->m_pointers[space]) { - if (pointer_record->m_owned[space]) { - void* space_ptr = pointer_record->m_pointers[space]; + if (space == spaceToFree || spaceToFree == NONE) { + if (pointer_record->m_pointers[space]) { + if (pointer_record->m_owned[space]) { + void* space_ptr = pointer_record->m_pointers[space]; #if defined(CHAI_ENABLE_UM) - if (space_ptr == pointer_record->m_pointers[UM]) { - callback(pointer_record, - ACTION_FREE, - ExecutionSpace(UM), - pointer_record->m_size); - { - std::lock_guard lock(m_mutex); - m_pointer_map.erase(space_ptr); - } - - auto alloc = m_resource_manager.getAllocator( - pointer_record->m_allocators[space]); - alloc.deallocate(space_ptr); - - for (int space_t = CPU; space_t < NUM_EXECUTION_SPACES; ++space_t) { - if (space_ptr == pointer_record->m_pointers[space_t]) - pointer_record->m_pointers[space_t] = nullptr; - } - } else { + if (space_ptr == pointer_record->m_pointers[UM]) { + callback(pointer_record, + ACTION_FREE, + ExecutionSpace(UM), + pointer_record->m_size); + { + std::lock_guard lock(m_mutex); + m_pointer_map.erase(space_ptr); + } + + auto alloc = m_resource_manager.getAllocator( + pointer_record->m_allocators[space]); + alloc.deallocate(space_ptr); + + for (int space_t = CPU; space_t < NUM_EXECUTION_SPACES; ++space_t) { + if (space_ptr == pointer_record->m_pointers[space_t]) + pointer_record->m_pointers[space_t] = nullptr; + } + } else { #endif - callback(pointer_record, - ACTION_FREE, - ExecutionSpace(space), - pointer_record->m_size); - { - std::lock_guard lock(m_mutex); - m_pointer_map.erase(space_ptr); - } - - auto alloc = m_resource_manager.getAllocator( - pointer_record->m_allocators[space]); - alloc.deallocate(space_ptr); - - pointer_record->m_pointers[space] = nullptr; + callback(pointer_record, + ACTION_FREE, + ExecutionSpace(space), + pointer_record->m_size); + { + std::lock_guard lock(m_mutex); + m_pointer_map.erase(space_ptr); + } + + auto alloc = m_resource_manager.getAllocator( + pointer_record->m_allocators[space]); + alloc.deallocate(space_ptr); + + pointer_record->m_pointers[space] = nullptr; #if defined(CHAI_ENABLE_UM) - } + } #endif + } } else { @@ -242,11 +244,12 @@ void ArrayManager::free(PointerRecord* pointer_record) } } } - - delete pointer_record; + + if (pointer_record != &s_null_record && spaceToFree == NONE) { + delete pointer_record; + } } - size_t ArrayManager::getSize(void* ptr) { // TODO @@ -370,4 +373,51 @@ ArrayManager::getAllocatorId(ExecutionSpace space) const } +void ArrayManager::evict(ExecutionSpace space, ExecutionSpace destinationSpace) { + // Check arguments + if (space == NONE) { + // Nothing to be done + return; + } + + if (destinationSpace == NONE) { + // If the destination space is NONE, evicting invalidates all data and + // leaves us in a bad state (if the last touch was in the eviction space). + CHAI_LOG(Warning, "evict does nothing with destinationSpace == NONE!"); + return; + } + + if (space == destinationSpace) { + // It doesn't make sense to evict to the same space, so do nothing + CHAI_LOG(Warning, "evict does nothing with space == destinationSpace!"); + return; + } + + // Now move and evict + std::vector pointersToEvict; + + for (auto entry : m_pointer_map) { + // Get the pointer record + auto record = *entry.second; + + // Move the data and register the touches + move(record, destinationSpace); + registerTouch(record, destinationSpace); + + // If the destinationSpace is ever allowed to be NONE, then we will need to + // update the touch in the eviction space and make sure the last space is not + // the eviction space. + + // Mark record for eviction later in this routine + pointersToEvict.push_back(record); + } + + // This must be done in a second pass because free erases from m_pointer_map, + // which would invalidate the iterator in the above loop + for (auto entry : pointersToEvict) { + free(entry, space); + } +} + + } // end of namespace chai diff --git a/src/chai/ArrayManager.hpp b/src/chai/ArrayManager.hpp index 9f4d52a1..6dc2e59f 100644 --- a/src/chai/ArrayManager.hpp +++ b/src/chai/ArrayManager.hpp @@ -147,9 +147,11 @@ class ArrayManager ExecutionSpace getDefaultAllocationSpace(); /*! - * \brief Free all allocations associated with the given PointerRecord. + * \brief Free allocation(s) associated with the given PointerRecord. + * Default (space == NONE) will free all allocations and delete + * the pointer record. */ - void free(PointerRecord* pointer); + void free(PointerRecord* pointer, ExecutionSpace space = NONE); #if defined(CHAI_ENABLE_PICK) template @@ -252,6 +254,15 @@ class ArrayManager */ bool deviceSynchronize() { return m_device_synchronize; } + /*! + * \brief Evicts the data in the given space. + * + * \param space Execution space to evict. + * \param destinationSpace The execution space to move the data to. + * Must not equal space or NONE. + */ + void evict(ExecutionSpace space, ExecutionSpace destinationSpace); + protected: /*! * \brief Construct a new ArrayManager. From de06319fcc73f2fafe4d191cc97b0493226635d4 Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Thu, 2 Apr 2020 15:35:30 -0700 Subject: [PATCH 10/70] build fixes --- examples/CMakeLists.txt | 7 ------- src/chai/ArrayManager.inl | 4 ++-- 2 files changed, 2 insertions(+), 9 deletions(-) diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt index 18e04049..e380076e 100644 --- a/examples/CMakeLists.txt +++ b/examples/CMakeLists.txt @@ -23,10 +23,6 @@ blt_add_executable( SOURCES chai-umpire-allocators.cpp DEPENDS_ON ${chai_umpire_example_depends}) -blt_add_test( - NAME managed_array_test - COMMAND managed_array_tests) - if (ENABLE_CUDA OR ENABLE_HIP) blt_add_executable( NAME chai-example.exe @@ -34,6 +30,3 @@ if (ENABLE_CUDA OR ENABLE_HIP) DEPENDS_ON ${chai_umpire_example_depends}) endif () -# if (ENABLE_RAJA_PLUGIN) -# add_subdirectory(integration) -# endif () diff --git a/src/chai/ArrayManager.inl b/src/chai/ArrayManager.inl index 6a56ac82..f8e72a1e 100644 --- a/src/chai/ArrayManager.inl +++ b/src/chai/ArrayManager.inl @@ -73,7 +73,7 @@ CHAI_INLINE typename ArrayManager::T_non_const ArrayManager::pick(T* src_ptr, size_t index) { T_non_const val; - m_resource_manager.registerAllocation(const_cast*>(&val), new umpire::util::AllocationRecord{const_cast*>(&val), sizeof(T), m_resource_manager.getAllocator("HOST").getAllocationStrategy()}); + m_resource_manager.registerAllocation(const_cast*>(&val), umpire::util::AllocationRecord{const_cast*>(&val), sizeof(T), m_resource_manager.getAllocator("HOST").getAllocationStrategy()}); m_resource_manager.copy(const_cast*>(&val), const_cast*>(src_ptr+index), sizeof(T)); m_resource_manager.deregisterAllocation(&val); return val; @@ -83,7 +83,7 @@ template CHAI_INLINE void ArrayManager::set(T* dst_ptr, size_t index, const T& val) { - m_resource_manager.registerAllocation(const_cast*>(&val), new umpire::util::AllocationRecord{const_cast*>(&val), sizeof(T), m_resource_manager.getAllocator("HOST").getAllocationStrategy()}); + m_resource_manager.registerAllocation(const_cast*>(&val), umpire::util::AllocationRecord{const_cast*>(&val), sizeof(T), m_resource_manager.getAllocator("HOST").getAllocationStrategy()}); m_resource_manager.copy(const_cast*>(dst_ptr+index), const_cast*>(&val), sizeof(T)); m_resource_manager.deregisterAllocation(const_cast*>(&val)); } From daee85de5cb29b262008ab2f1e198add9824f5bf Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Thu, 2 Apr 2020 16:06:39 -0700 Subject: [PATCH 11/70] fix build issue in managed_array_test --- src/chai/ManagedArray_thin.inl | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/chai/ManagedArray_thin.inl b/src/chai/ManagedArray_thin.inl index 83bac160..a664793b 100644 --- a/src/chai/ManagedArray_thin.inl +++ b/src/chai/ManagedArray_thin.inl @@ -158,7 +158,7 @@ CHAI_HOST void ManagedArray::free() #if defined(CHAI_ENABLE_UM) cudaFree(m_active_base_pointer); #else - ::free(m_active_base_pointer); + ::free((void *)m_active_base_pointer); #endif m_active_base_pointer = nullptr; From 00f48a25a898def15053c2f18b7c1e4a4e7bca26 Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Fri, 3 Apr 2020 06:53:30 -0700 Subject: [PATCH 12/70] add ouput-on-failure option to ctest calls. --- Dockerfile | 14 +++++++------- 1 file changed, 7 insertions(+), 7 deletions(-) diff --git a/Dockerfile b/Dockerfile index 5b49f084..31cb0eed 100644 --- a/Dockerfile +++ b/Dockerfile @@ -3,49 +3,49 @@ COPY --chown=axom:axom . /home/axom/workspace WORKDIR /home/axom/workspace RUN mkdir build && cd build && cmake -DCMAKE_CXX_COMPILER=g++ -DENABLE_CUDA=OFF .. RUN cd build && make -j 16 -RUN cd build && ctest -T test +RUN cd build && ctest -T test --output-on-failure FROM axom/compilers:gcc-6 AS gcc6 COPY --chown=axom:axom . /home/axom/workspace WORKDIR /home/axom/workspace RUN mkdir build && cd build && cmake -DCMAKE_CXX_COMPILER=g++ -DENABLE_CUDA=OFF .. RUN cd build && make -j 16 -RUN cd build && ctest -T test +RUN cd build && ctest -T test --output-on-failure FROM axom/compilers:gcc-7 AS gcc7 COPY --chown=axom:axom . /home/axom/workspace WORKDIR /home/axom/workspace RUN mkdir build && cd build && cmake -DCMAKE_CXX_COMPILER=g++ -DENABLE_CUDA=OFF .. RUN cd build && make -j 16 -RUN cd build && ctest -T test +RUN cd build && ctest -T test --output-on-failure FROM axom/compilers:gcc-8 AS gcc8 COPY --chown=axom:axom . /home/axom/workspace WORKDIR /home/axom/workspace RUN mkdir build && cd build && cmake -DCMAKE_CXX_COMPILER=g++ -DENABLE_CUDA=OFF .. RUN cd build && make -j 16 -RUN cd build && ctest -T test +RUN cd build && ctest -T test --output-on-failure FROM axom/compilers:clang-4 AS clang4 COPY --chown=axom:axom . /home/axom/workspace WORKDIR /home/axom/workspace RUN mkdir build && cd build && cmake -DCMAKE_CXX_COMPILER=clang++ -DENABLE_CUDA=OFF .. RUN cd build && make -j 16 -RUN cd build && ctest -T test +RUN cd build && ctest -T test --output-on-failure FROM axom/compilers:clang-5 AS clang5 COPY --chown=axom:axom . /home/axom/workspace WORKDIR /home/axom/workspace RUN mkdir build && cd build && cmake -DCMAKE_CXX_COMPILER=clang++ -DENABLE_CUDA=OFF .. RUN cd build && make -j 16 -RUN cd build && ctest -T test +RUN cd build && ctest -T test --output-on-failure FROM axom/compilers:clang-6 AS clang6 COPY --chown=axom:axom . /home/axom/workspace WORKDIR /home/axom/workspace RUN mkdir build && cd build && cmake -DCMAKE_CXX_COMPILER=clang++ -DENABLE_CUDA=OFF .. RUN cd build && make -j 16 -RUN cd build && ctest -T test +RUN cd build && ctest -T test --output-on-failure FROM axom/compilers:nvcc-9 AS nvcc From e616f6e8e37c32e9924c99decd3049ea1eb60bec Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Fri, 3 Apr 2020 07:21:46 -0700 Subject: [PATCH 13/70] removing null_record protection as an experiment. --- src/chai/ArrayManager.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/chai/ArrayManager.cpp b/src/chai/ArrayManager.cpp index f193a05d..9319841b 100644 --- a/src/chai/ArrayManager.cpp +++ b/src/chai/ArrayManager.cpp @@ -245,7 +245,7 @@ void ArrayManager::free(PointerRecord* pointer_record, ExecutionSpace spaceToFre } } - if (pointer_record != &s_null_record && spaceToFree == NONE) { + if (spaceToFree == NONE) { delete pointer_record; } } From 556cbc0054d132d786170212994effe4147362b9 Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Fri, 3 Apr 2020 07:28:32 -0700 Subject: [PATCH 14/70] Revert "removing null_record protection as an experiment." This reverts commit e616f6e8e37c32e9924c99decd3049ea1eb60bec. --- src/chai/ArrayManager.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/chai/ArrayManager.cpp b/src/chai/ArrayManager.cpp index 9319841b..f193a05d 100644 --- a/src/chai/ArrayManager.cpp +++ b/src/chai/ArrayManager.cpp @@ -245,7 +245,7 @@ void ArrayManager::free(PointerRecord* pointer_record, ExecutionSpace spaceToFre } } - if (spaceToFree == NONE) { + if (pointer_record != &s_null_record && spaceToFree == NONE) { delete pointer_record; } } From 83a4f63fa4f21689c7e40bde050c7355bd288701 Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Fri, 3 Apr 2020 07:34:48 -0700 Subject: [PATCH 15/70] restore duplicate instance of managed_array_test.. --- examples/CMakeLists.txt | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt index e380076e..d029707a 100644 --- a/examples/CMakeLists.txt +++ b/examples/CMakeLists.txt @@ -23,6 +23,10 @@ blt_add_executable( SOURCES chai-umpire-allocators.cpp DEPENDS_ON ${chai_umpire_example_depends}) +blt_add_test( + NAME managed_array_test + COMMAND managed_array_tests) + if (ENABLE_CUDA OR ENABLE_HIP) blt_add_executable( NAME chai-example.exe From 3c003220d4ea2931e0e58b682b4b1366ae4349db Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Fri, 3 Apr 2020 08:00:31 -0700 Subject: [PATCH 16/70] to the printf debugging machine! --- src/chai/ArrayManager.cpp | 1 + tests/integration/managed_array_tests.cpp | 4 ++++ 2 files changed, 5 insertions(+) diff --git a/src/chai/ArrayManager.cpp b/src/chai/ArrayManager.cpp index f193a05d..8af2e58a 100644 --- a/src/chai/ArrayManager.cpp +++ b/src/chai/ArrayManager.cpp @@ -198,6 +198,7 @@ void ArrayManager::free(PointerRecord* pointer_record, ExecutionSpace spaceToFre if (pointer_record->m_pointers[space]) { if (pointer_record->m_owned[space]) { void* space_ptr = pointer_record->m_pointers[space]; + printf("in free with space_ptr %p %b \n",space_ptr, pointer_record==&s_null_record); #if defined(CHAI_ENABLE_UM) if (space_ptr == pointer_record->m_pointers[UM]) { callback(pointer_record, diff --git a/tests/integration/managed_array_tests.cpp b/tests/integration/managed_array_tests.cpp index 95d40bea..411929f2 100644 --- a/tests/integration/managed_array_tests.cpp +++ b/tests/integration/managed_array_tests.cpp @@ -631,11 +631,15 @@ GPU_TEST(ManagedArray, ReallocateGPU) TEST(ManagedArray, NullpointerConversions) { chai::ManagedArray a; + printf("a.free()\n"); a.free(); + printf("a to nullptr\n"); a = nullptr; chai::ManagedArray b; + printf("b.free()\n"); b.free(); + printf("b to nullptr\n"); b = nullptr; ASSERT_EQ(a.size(), 0u); From c6f053b6652f45bc065ed6be9238aae80fdb6d9a Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Fri, 3 Apr 2020 08:00:57 -0700 Subject: [PATCH 17/70] to the printf debugging machine! --- tests/integration/managed_array_tests.cpp | 1 + 1 file changed, 1 insertion(+) diff --git a/tests/integration/managed_array_tests.cpp b/tests/integration/managed_array_tests.cpp index 411929f2..e86f3241 100644 --- a/tests/integration/managed_array_tests.cpp +++ b/tests/integration/managed_array_tests.cpp @@ -645,6 +645,7 @@ TEST(ManagedArray, NullpointerConversions) ASSERT_EQ(a.size(), 0u); ASSERT_EQ(b.size(), 0u); + printf("c from nullptr\n"); chai::ManagedArray c(nullptr); ASSERT_EQ(c.size(), 0u); From 3bf96905c47e6e30c06d028d75ac32e10ef22ad8 Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Sat, 4 Apr 2020 08:00:48 -0700 Subject: [PATCH 18/70] fix test failure, clean up debugging code. --- examples/CMakeLists.txt | 4 ---- src/chai/ArrayManager.cpp | 9 ++++----- tests/integration/managed_array_tests.cpp | 5 ----- 3 files changed, 4 insertions(+), 14 deletions(-) diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt index d029707a..e380076e 100644 --- a/examples/CMakeLists.txt +++ b/examples/CMakeLists.txt @@ -23,10 +23,6 @@ blt_add_executable( SOURCES chai-umpire-allocators.cpp DEPENDS_ON ${chai_umpire_example_depends}) -blt_add_test( - NAME managed_array_test - COMMAND managed_array_tests) - if (ENABLE_CUDA OR ENABLE_HIP) blt_add_executable( NAME chai-example.exe diff --git a/src/chai/ArrayManager.cpp b/src/chai/ArrayManager.cpp index 8af2e58a..1c96706c 100644 --- a/src/chai/ArrayManager.cpp +++ b/src/chai/ArrayManager.cpp @@ -198,7 +198,6 @@ void ArrayManager::free(PointerRecord* pointer_record, ExecutionSpace spaceToFre if (pointer_record->m_pointers[space]) { if (pointer_record->m_owned[space]) { void* space_ptr = pointer_record->m_pointers[space]; - printf("in free with space_ptr %p %b \n",space_ptr, pointer_record==&s_null_record); #if defined(CHAI_ENABLE_UM) if (space_ptr == pointer_record->m_pointers[UM]) { callback(pointer_record, @@ -238,10 +237,10 @@ void ArrayManager::free(PointerRecord* pointer_record, ExecutionSpace spaceToFre } #endif } - } - else - { - m_resource_manager.deregisterAllocation(pointer_record->m_pointers[space]); + else + { + m_resource_manager.deregisterAllocation(pointer_record->m_pointers[space]); + } } } } diff --git a/tests/integration/managed_array_tests.cpp b/tests/integration/managed_array_tests.cpp index e86f3241..95d40bea 100644 --- a/tests/integration/managed_array_tests.cpp +++ b/tests/integration/managed_array_tests.cpp @@ -631,21 +631,16 @@ GPU_TEST(ManagedArray, ReallocateGPU) TEST(ManagedArray, NullpointerConversions) { chai::ManagedArray a; - printf("a.free()\n"); a.free(); - printf("a to nullptr\n"); a = nullptr; chai::ManagedArray b; - printf("b.free()\n"); b.free(); - printf("b to nullptr\n"); b = nullptr; ASSERT_EQ(a.size(), 0u); ASSERT_EQ(b.size(), 0u); - printf("c from nullptr\n"); chai::ManagedArray c(nullptr); ASSERT_EQ(c.size(), 0u); From 89d368bc167cde95c67eef09a64796cac25ff047 Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Mon, 6 Apr 2020 12:24:21 -0700 Subject: [PATCH 19/70] windows DLL build fixes. --- src/chai/ArrayManager.hpp | 47 ++++++++++++++------------- src/chai/RajaExecutionSpacePlugin.cpp | 15 +++++++++ src/chai/Types.hpp | 12 ++++++- 3 files changed, 50 insertions(+), 24 deletions(-) diff --git a/src/chai/ArrayManager.hpp b/src/chai/ArrayManager.hpp index 6dc2e59f..6b62d6df 100644 --- a/src/chai/ArrayManager.hpp +++ b/src/chai/ArrayManager.hpp @@ -52,7 +52,7 @@ class ArrayManager using PointerMap = umpire::util::MemoryMap; - static PointerRecord s_null_record; + CHAISHAREDDLL_API static PointerRecord s_null_record; /*! * \brief Get the singleton instance. @@ -60,6 +60,7 @@ class ArrayManager * \return Pointer to the ArrayManager instance. * */ + CHAISHAREDDLL_API CHAI_HOST_DEVICE static ArrayManager* getInstance(); @@ -68,14 +69,14 @@ class ArrayManager * * \param space The space to set as current. */ - void setExecutionSpace(ExecutionSpace space); + CHAISHAREDDLL_API void setExecutionSpace(ExecutionSpace space); /*! * \brief Get the current execution space. * * \return The current execution space.jo */ - ExecutionSpace getExecutionSpace(); + CHAISHAREDDLL_API ExecutionSpace getExecutionSpace(); /*! * \brief Move data in pointer to the current execution space. @@ -83,16 +84,16 @@ class ArrayManager * \param pointer Pointer to data in any execution space. * \return Pointer to data in the current execution space. */ - void* move(void* pointer, - PointerRecord* pointer_record, - ExecutionSpace = NONE); + CHAISHAREDDLL_API void* move(void* pointer, + PointerRecord* pointer_record, + ExecutionSpace = NONE); /*! * \brief Register a touch of the pointer in the current execution space. * * \param pointer Raw pointer to register a touch of. */ - void registerTouch(PointerRecord* pointer_record); + CHAISHAREDDLL_API void registerTouch(PointerRecord* pointer_record); /*! * \brief Register a touch of the pointer in the given execution space. @@ -102,7 +103,7 @@ class ArrayManager * \param pointer Raw pointer to register a touch of. * \param space Space to register touch. */ - void registerTouch(PointerRecord* pointer_record, ExecutionSpace space); + CHAISHAREDDLL_API void registerTouch(PointerRecord* pointer_record, ExecutionSpace space); /*! * \brief Make a new allocation of the data described by the PointerRecord in @@ -111,7 +112,7 @@ class ArrayManager * \param pointer_record * \param space Space in which to make the allocation. */ - void allocate(PointerRecord* pointer_record, ExecutionSpace space = CPU); + CHAISHAREDDLL_API void allocate(PointerRecord* pointer_record, ExecutionSpace space = CPU); /*! * \brief Reallocate data. @@ -135,7 +136,7 @@ class ArrayManager * * \param space New space for default allocations. */ - void setDefaultAllocationSpace(ExecutionSpace space); + CHAISHAREDDLL_API void setDefaultAllocationSpace(ExecutionSpace space); /*! * \brief Get the currently set default allocation space. @@ -144,14 +145,14 @@ class ArrayManager * * \return Current default space for allocations. */ - ExecutionSpace getDefaultAllocationSpace(); + CHAISHAREDDLL_API ExecutionSpace getDefaultAllocationSpace(); /*! * \brief Free allocation(s) associated with the given PointerRecord. * Default (space == NONE) will free all allocations and delete * the pointer record. */ - void free(PointerRecord* pointer, ExecutionSpace space = NONE); + CHAISHAREDDLL_API void free(PointerRecord* pointer, ExecutionSpace space = NONE); #if defined(CHAI_ENABLE_PICK) template @@ -167,9 +168,9 @@ class ArrayManager * \param pointer Pointer to find the size of. * \return Size of pointer. */ - size_t getSize(void* pointer); + CHAISHAREDDLL_API size_t getSize(void* pointer); - PointerRecord* makeManaged(void* pointer, + CHAISHAREDDLL_API PointerRecord* makeManaged(void* pointer, size_t size, ExecutionSpace space, bool owned); @@ -177,14 +178,14 @@ class ArrayManager /*! * \brief Assign a user-defined callback triggerd upon memory operations. */ - void setUserCallback(void* pointer, UserCallback const& f); + CHAISHAREDDLL_API void setUserCallback(void* pointer, UserCallback const& f); /*! * \brief Set touched to false in all spaces for the given PointerRecord. * * \param pointer_record PointerRecord to reset. */ - void resetTouch(PointerRecord* pointer_record); + CHAISHAREDDLL_API void resetTouch(PointerRecord* pointer_record); /*! * \brief Find the PointerRecord corresponding to the raw pointer. @@ -194,7 +195,7 @@ class ArrayManager * \return PointerRecord containing the raw pointer, or an empty * PointerRecord if none found. */ - PointerRecord* getPointerRecord(void* pointer); + CHAISHAREDDLL_API PointerRecord* getPointerRecord(void* pointer); /*! * \brief Create a copy of the given PointerRecord with a new allocation @@ -204,30 +205,30 @@ class ArrayManager * * \return A copy of the given PointerRecord, must be free'd with delete. */ - PointerRecord* deepCopyRecord(PointerRecord const* record); + CHAISHAREDDLL_API PointerRecord* deepCopyRecord(PointerRecord const* record); /*! * \brief Create a copy of the pointer map. * * \return A copy of the pointer map. Can be used to find memory leaks. */ - std::unordered_map getPointerMap() const; + CHAISHAREDDLL_API std::unordered_map getPointerMap() const; /*! * \brief Get the total number of arrays registered with the array manager. * * \return The total number of arrays registered with the array manager. */ - size_t getTotalNumArrays() const; + CHAISHAREDDLL_API size_t getTotalNumArrays() const; /*! * \brief Get the total amount of memory allocated. * * \return The total amount of memory allocated. */ - size_t getTotalSize() const; + CHAISHAREDDLL_API size_t getTotalSize() const; - int getAllocatorId(ExecutionSpace space) const; + CHAISHAREDDLL_API int getAllocatorId(ExecutionSpace space) const; /*! * \brief Turn callbacks on. @@ -261,7 +262,7 @@ class ArrayManager * \param destinationSpace The execution space to move the data to. * Must not equal space or NONE. */ - void evict(ExecutionSpace space, ExecutionSpace destinationSpace); + CHAISHAREDDLL_API void evict(ExecutionSpace space, ExecutionSpace destinationSpace); protected: /*! diff --git a/src/chai/RajaExecutionSpacePlugin.cpp b/src/chai/RajaExecutionSpacePlugin.cpp index 71fe7716..6ba50463 100644 --- a/src/chai/RajaExecutionSpacePlugin.cpp +++ b/src/chai/RajaExecutionSpacePlugin.cpp @@ -76,6 +76,21 @@ RajaExecutionSpacePlugin::postLaunch(RAJA::util::PluginContext) } +// 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 P( "RajaExecutionSpacePlugin", diff --git a/src/chai/Types.hpp b/src/chai/Types.hpp index 37a57b86..48a968ad 100644 --- a/src/chai/Types.hpp +++ b/src/chai/Types.hpp @@ -9,6 +9,16 @@ #include +#if defined(_WIN32) && !defined(CHAISTATICLIB) +#ifdef CHAISHAREDDLL_EXPORTS +#define CHAISHAREDDLL_API __declspec(dllexport) +#else +#define CHAISHAREDDLL_API __declspec(dllimport) +#endif +#else +#define CHAISHAREDDLL_API +#endif + namespace chai { @@ -16,7 +26,7 @@ typedef unsigned int uint; enum Action { ACTION_ALLOC, ACTION_FREE, ACTION_MOVE }; -using UserCallback = std::function; +using UserCallback = std::function; } // end of namespace chai From 892f05a340b074ca5166b55bfec54b2791918f15 Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Tue, 7 Apr 2020 07:53:48 -0700 Subject: [PATCH 20/70] support device side slicing, fix bug where T* return wrong pointer for slices, move slice definitions to common implementation in ManagedArray.hpp --- src/chai/ManagedArray.hpp | 44 ++++++++++++++++++++++++++++------ src/chai/ManagedArray.inl | 20 +--------------- src/chai/ManagedArray_thin.inl | 13 ---------- 3 files changed, 38 insertions(+), 39 deletions(-) diff --git a/src/chai/ManagedArray.hpp b/src/chai/ManagedArray.hpp index ab30a868..a4513516 100644 --- a/src/chai/ManagedArray.hpp +++ b/src/chai/ManagedArray.hpp @@ -156,7 +156,7 @@ class ManagedArray : public CHAICopyable CHAI_HOST void move(ExecutionSpace space); - CHAI_HOST ManagedArray slice(size_t begin, size_t end); + CHAI_HOST_DEVICE ManagedArray slice(size_t begin, size_t elems=-1) const; /*! * \brief Return reference to i-th element of the ManagedArray. * @@ -171,7 +171,13 @@ class ManagedArray : public CHAICopyable * \brief get access to m_active_pointer * @return a copy of m_active_pointer */ - T* getActiveBasePointer() const; + CHAI_HOST_DEVICE T* getActiveBasePointer() const; + + /*! + * \brief get access to m_active_pointer + * @return a copy of m_active_pointer + */ + CHAI_HOST_DEVICE T* getActivePointer() const; /*! * \brief @@ -314,20 +320,20 @@ class ManagedArray : public CHAICopyable /*! * Pointer to ArrayManager instance. */ - ArrayManager* m_resource_manager = nullptr; + mutable ArrayManager* m_resource_manager = nullptr; /*! * Number of elements in the ManagedArray. */ - size_t m_elems = 0; - size_t m_offset = 0; + mutable size_t m_elems = 0; + mutable size_t m_offset = 0; /*! * Pointer to PointerRecord data. */ - PointerRecord* m_pointer_record = nullptr; + mutable PointerRecord* m_pointer_record = nullptr; - bool m_is_slice = false; + mutable bool m_is_slice = false; }; @@ -396,6 +402,30 @@ ManagedArray deepCopy(ManagedArray const& array) return ManagedArray(copy_record, copy_record->m_last_space); } +template +CHAI_INLINE CHAI_HOST_DEVICE ManagedArray ManagedArray::slice( size_t offset, size_t elems) const +{ + ManagedArray slice; + slice.m_resource_manager = m_resource_manager; + if (elems == -1) { + elems = size() - offset; + } + if (offset + elems > size()) { +#ifndef __CUDA_ARCH__ + CHAI_LOG(Debug, + "Invalid slice. No active pointer or index out of bounds"); +#endif + } else { + slice.m_pointer_record = m_pointer_record; + slice.m_active_base_pointer = m_active_base_pointer; + slice.m_offset = offset + m_offset; + slice.m_active_pointer = m_active_base_pointer + slice.m_offset; + slice.m_elems = elems; + slice.m_is_slice = true; + } + return slice; +} + } // end of namespace chai #if defined(CHAI_DISABLE_RM) diff --git a/src/chai/ManagedArray.inl b/src/chai/ManagedArray.inl index d3b809d5..6279b375 100644 --- a/src/chai/ManagedArray.inl +++ b/src/chai/ManagedArray.inl @@ -155,24 +155,6 @@ CHAI_HOST_DEVICE ManagedArray::ManagedArray(T* data, ArrayManager* array_mana { } -template -CHAI_INLINE -CHAI_HOST ManagedArray ManagedArray::slice(size_t offset, size_t elems) { - ManagedArray slice(nullptr); - slice.m_resource_manager = m_resource_manager; - if(offset + elems > size()) { - CHAI_LOG(Debug, "Invalid slice. No active pointer or index out of bounds"); - } else { - slice.m_pointer_record = m_pointer_record; - slice.m_active_base_pointer = m_active_base_pointer; - slice.m_offset = offset + m_offset; - slice.m_active_pointer = m_active_base_pointer + slice.m_offset; - slice.m_elems = elems; - slice.m_is_slice = true; - } - return slice; -} - template CHAI_HOST void ManagedArray::allocate( size_t elems, @@ -365,7 +347,7 @@ CHAI_HOST_DEVICE ManagedArray::operator T*() const { m_resource_manager->setExecutionSpace(CPU); auto non_const_active_base_pointer = const_cast(static_cast(m_active_base_pointer)); m_active_base_pointer = static_cast(m_resource_manager->move(non_const_active_base_pointer, m_pointer_record)); - m_active_pointer = m_active_base_pointer; + m_active_pointer = m_active_base_pointer + m_offset; m_resource_manager->registerTouch(m_pointer_record); diff --git a/src/chai/ManagedArray_thin.inl b/src/chai/ManagedArray_thin.inl index a664793b..014cd286 100644 --- a/src/chai/ManagedArray_thin.inl +++ b/src/chai/ManagedArray_thin.inl @@ -80,19 +80,6 @@ CHAI_HOST_DEVICE ManagedArray::ManagedArray(T* data, { } -template -CHAI_INLINE -CHAI_HOST ManagedArray ManagedArray::slice(size_t offset, size_t elems) { - ManagedArray slice; - if (offset + elems > size()) { - CHAI_LOG(Debug, "Invalid slice. No active pointer or index out of bounds"); - } else { - slice.m_active_pointer = m_active_pointer + offset; - slice.m_elems = elems; - slice.m_is_slice = true; - } - return slice; -} template CHAI_INLINE From d936314d3cb3b0b778627925c3e62db5f99a3269 Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Tue, 7 Apr 2020 08:20:03 -0700 Subject: [PATCH 21/70] add missing getActivePointer definitions --- src/chai/ManagedArray.inl | 6 ++++++ src/chai/ManagedArray_thin.inl | 11 +++++++++++ 2 files changed, 17 insertions(+) diff --git a/src/chai/ManagedArray.inl b/src/chai/ManagedArray.inl index 6279b375..69a9708d 100644 --- a/src/chai/ManagedArray.inl +++ b/src/chai/ManagedArray.inl @@ -390,6 +390,12 @@ ManagedArray::getActiveBasePointer() const return m_active_base_pointer; } +template +T* +ManagedArray::getActivePointer() const +{ + return m_active_pointer; +} //template //ManagedArray::operator ManagedArray< diff --git a/src/chai/ManagedArray_thin.inl b/src/chai/ManagedArray_thin.inl index 014cd286..db06eaf4 100644 --- a/src/chai/ManagedArray_thin.inl +++ b/src/chai/ManagedArray_thin.inl @@ -80,6 +80,17 @@ CHAI_HOST_DEVICE ManagedArray::ManagedArray(T* data, { } +template +T* ManagedArray::getActiveBasePointer() const +{ + return m_active_base_pointer; +} + +template +T* ManagedArray::getActivePointer() const +{ + return m_active_pointer; +} template CHAI_INLINE From d402aaab320133819f660a2e0b82fe75296f3fb3 Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Tue, 7 Apr 2020 09:17:14 -0700 Subject: [PATCH 22/70] DEBUG->CHAI_DEBUG, add a getPointer(space), get more accurate m_elems count --- CMakeLists.txt | 1 + src/chai/ArrayManager.inl | 1 + src/chai/ChaiMacros.hpp | 2 +- src/chai/ManagedArray.hpp | 9 +++++++++ src/chai/ManagedArray.inl | 21 ++++++++++++++++++++- src/chai/ManagedArray_thin.inl | 6 ++++++ src/chai/config.hpp.in | 1 + 7 files changed, 39 insertions(+), 2 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 6731b652..8fa9af23 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -21,6 +21,7 @@ mark_as_advanced(DISABLE_RM) option(ENABLE_UM "Use CUDA unified (managed) memory" Off) option(ENABLE_RAJA_PLUGIN "Build plugin to set RAJA execution spaces" Off) option(CHAI_ENABLE_GPU_ERROR_CHECKING "Enable GPU error checking" On) +option(CHAI_DEBUG "Enable Debug Logging.") set(ENABLE_TESTS On CACHE BOOL "") set(ENABLE_EXAMPLES On CACHE BOOL "") diff --git a/src/chai/ArrayManager.inl b/src/chai/ArrayManager.inl index f8e72a1e..0ccc35fd 100644 --- a/src/chai/ArrayManager.inl +++ b/src/chai/ArrayManager.inl @@ -31,6 +31,7 @@ void* ArrayManager::reallocate(void* pointer, size_t elems, PointerRecord* point for (int space = CPU; space < NUM_EXECUTION_SPACES; ++space) { if (pointer_record->m_pointers[space] == pointer) { my_space = static_cast(space); + break; } } diff --git a/src/chai/ChaiMacros.hpp b/src/chai/ChaiMacros.hpp index 295c4856..6fd48d7d 100644 --- a/src/chai/ChaiMacros.hpp +++ b/src/chai/ChaiMacros.hpp @@ -44,7 +44,7 @@ #else -#if defined(DEBUG) +#if defined(CHAI_DEBUG) #define CHAI_LOG(level, msg) \ std::cerr << "[" << __FILE__ << "] " << msg << std::endl; diff --git a/src/chai/ManagedArray.hpp b/src/chai/ManagedArray.hpp index a4513516..63996c61 100644 --- a/src/chai/ManagedArray.hpp +++ b/src/chai/ManagedArray.hpp @@ -179,6 +179,15 @@ class ManagedArray : public CHAICopyable */ CHAI_HOST_DEVICE T* getActivePointer() const; + /*! + * \brief get access to the pointer in the given execution space + * @return a copy of the pointer in the given execution space + * + * \param space The space to get the pointer for. + * \param do_move Ensure data at that pointer is live and valid. + */ + CHAI_HOST T* getPointer(ExecutionSpace space, bool do_move = true); + /*! * \brief * diff --git a/src/chai/ManagedArray.inl b/src/chai/ManagedArray.inl index 69a9708d..89fe959c 100644 --- a/src/chai/ManagedArray.inl +++ b/src/chai/ManagedArray.inl @@ -138,6 +138,9 @@ CHAI_HOST_DEVICE ManagedArray::ManagedArray(ManagedArray const& other): m_is_slice(other.m_is_slice) { #if !defined(__CUDA_ARCH__) && !defined(__HIP_DEVICE_COMPILE__) + if (other.m_pointer_record != nullptr) { + m_elems = other.m_pointer_record->m_size/sizeof(T); + } move(m_resource_manager->getExecutionSpace()); #endif } @@ -370,7 +373,7 @@ CHAI_HOST_DEVICE ManagedArray::ManagedArray(T* data, bool ) : m_active_base_pointer(data), #if !defined(__CUDA_ARCH__) && !defined(__HIP_DEVICE_COMPILE__) m_resource_manager(ArrayManager::getInstance()), - m_elems(m_resource_manager->getSize(m_active_base_pointer)), + m_elems(m_resource_manager->getSize((void *)m_active_base_pointer)/sizeof(T)), m_pointer_record(m_resource_manager->getPointerRecord(data)), #else m_resource_manager(nullptr), @@ -397,6 +400,22 @@ ManagedArray::getActivePointer() const return m_active_pointer; } +template +T* +ManagedArray::getPointer(ExecutionSpace space, bool do_move) { + if (m_elems == 0 && !m_is_slice) { + return nullptr; + } + if (do_move) { + ExecutionSpace oldContext = m_resource_manager->getExecutionSpace(); + m_resource_manager->setExecutionSpace(space); + move(space); + m_resource_manager->setExecutionSpace(oldContext); + } + int offset = m_is_slice ? m_offset : 0 ; + return ((T*) m_pointer_record->m_pointers[space]) + offset; +} + //template //ManagedArray::operator ManagedArray< // typename std::conditional::value, diff --git a/src/chai/ManagedArray_thin.inl b/src/chai/ManagedArray_thin.inl index db06eaf4..80f52549 100644 --- a/src/chai/ManagedArray_thin.inl +++ b/src/chai/ManagedArray_thin.inl @@ -92,6 +92,12 @@ T* ManagedArray::getActivePointer() const return m_active_pointer; } +template +T* ManagedArray::getPointer(ExecutionSpace /*space*/, bool) +{ + return m_active_pointer; +} + template CHAI_INLINE CHAI_HOST void ManagedArray::allocate(size_t elems, diff --git a/src/chai/config.hpp.in b/src/chai/config.hpp.in index ae420f5e..5a1c53ac 100644 --- a/src/chai/config.hpp.in +++ b/src/chai/config.hpp.in @@ -13,6 +13,7 @@ #cmakedefine CHAI_ENABLE_IMPLICIT_CONVERSIONS #cmakedefine CHAI_DISABLE_RM #cmakedefine CHAI_ENABLE_UM +#cmakedefine CHAI_DEBUG #cmakedefine CHAI_ENABLE_RAJA_PLUGIN #cmakedefine CHAI_ENABLE_GPU_SIMULATION_MODE From d8c0ef7fa34bc71308e98000cc4c53d8d86e91df Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Tue, 7 Apr 2020 10:49:36 -0700 Subject: [PATCH 23/70] Add ArrayManager::copy and miscelaneous operators to ManagedArray --- src/chai/ArrayManager.hpp | 6 ++++ src/chai/ArrayManager.inl | 5 +++ src/chai/ManagedArray.hpp | 11 ++++++- src/chai/ManagedArray.inl | 57 +++++++++++++++++++++++++++++++--- src/chai/ManagedArray_thin.inl | 53 ++++++++++++++++++++++++++++--- src/chai/managed_ptr.hpp | 2 +- 6 files changed, 124 insertions(+), 10 deletions(-) diff --git a/src/chai/ArrayManager.hpp b/src/chai/ArrayManager.hpp index 6b62d6df..1279944b 100644 --- a/src/chai/ArrayManager.hpp +++ b/src/chai/ArrayManager.hpp @@ -230,6 +230,12 @@ class ArrayManager CHAISHAREDDLL_API int getAllocatorId(ExecutionSpace space) const; + /*! + * \brief Wraps our resource manager's copy. + * + */ + CHAISHAREDDLL_API void copy(void * dst, void * src, size_t size); + /*! * \brief Turn callbacks on. */ diff --git a/src/chai/ArrayManager.inl b/src/chai/ArrayManager.inl index 0ccc35fd..ab2e4475 100644 --- a/src/chai/ArrayManager.inl +++ b/src/chai/ArrayManager.inl @@ -90,6 +90,11 @@ void ArrayManager::set(T* dst_ptr, size_t index, const T& val) } #endif +CHAI_INLINE +void ArrayManager::copy(void * dst, void * src, size_t size) { + m_resource_manager.copy(dst,src,size); +} + CHAI_INLINE void ArrayManager::setAllocator(ExecutionSpace space, umpire::Allocator &allocator) { *m_allocators[space] = allocator; diff --git a/src/chai/ManagedArray.hpp b/src/chai/ManagedArray.hpp index 63996c61..72b04789 100644 --- a/src/chai/ManagedArray.hpp +++ b/src/chai/ManagedArray.hpp @@ -210,8 +210,17 @@ class ManagedArray : public CHAICopyable CHAI_HOST_DEVICE ManagedArray& operator=(std::nullptr_t); - CHAI_HOST_DEVICE bool operator==(ManagedArray& rhs); + CHAI_HOST_DEVICE bool operator==(ManagedArray& rhs) const; + CHAI_HOST_DEVICE bool operator!=(ManagedArray& from) const; + CHAI_HOST_DEVICE bool operator==(T* from) const; + CHAI_HOST_DEVICE bool operator!=(T* from) const; + + CHAI_HOST_DEVICE bool operator==(std::nullptr_t from) const; + CHAI_HOST_DEVICE bool operator!=(std::nullptr_t from) const; + + + CHAI_HOST_DEVICE explicit operator bool() const; #if defined(CHAI_ENABLE_PICK) /*! diff --git a/src/chai/ManagedArray.inl b/src/chai/ManagedArray.inl index 89fe959c..c16bb56f 100644 --- a/src/chai/ManagedArray.inl +++ b/src/chai/ManagedArray.inl @@ -400,10 +400,10 @@ ManagedArray::getActivePointer() const return m_active_pointer; } -template +template T* -ManagedArray::getPointer(ExecutionSpace space, bool do_move) { - if (m_elems == 0 && !m_is_slice) { +ManagedArray::getPointer(ExecutionSpace space, bool do_move) { + if (m_elems == 0 && !m_is_slice) { return nullptr; } if (do_move) { @@ -460,10 +460,59 @@ template CHAI_INLINE CHAI_HOST_DEVICE bool -ManagedArray::operator== (ManagedArray& rhs) { +ManagedArray::operator== (ManagedArray& rhs) const +{ return (m_active_pointer == rhs.m_active_pointer); } +template +CHAI_INLINE +CHAI_HOST_DEVICE +bool +ManagedArray::operator!= (ManagedArray& rhs) const +{ + return (m_active_pointer != rhs.m_active_pointer); +} + + +template +CHAI_INLINE +CHAI_HOST_DEVICE +bool +ManagedArray::operator== (T * from) const { + return m_active_pointer == from; +} + +template +CHAI_INLINE +CHAI_HOST_DEVICE +bool +ManagedArray::operator!= (T * from) const { + return m_active_pointer != from; +} + +template +CHAI_INLINE +CHAI_HOST_DEVICE +bool +ManagedArray::operator== (std::nullptr_t from) const { + return m_active_pointer == from || m_elems == 0; +} +template +CHAI_INLINE +CHAI_HOST_DEVICE +bool +ManagedArray::operator!= (std::nullptr_t from) const { + return m_active_pointer != from && m_elems > 0; +} + +template +CHAI_INLINE +CHAI_HOST_DEVICE +ManagedArray::operator bool () const { + return m_elems > 0; +} + template template::type> CHAI_INLINE diff --git a/src/chai/ManagedArray_thin.inl b/src/chai/ManagedArray_thin.inl index 80f52549..834eef3f 100644 --- a/src/chai/ManagedArray_thin.inl +++ b/src/chai/ManagedArray_thin.inl @@ -261,15 +261,60 @@ typename std::enable_if< !std::is_const::value , } template -CHAI_INLINE -CHAI_HOST_DEVICE -ManagedArray& -ManagedArray::operator= (std::nullptr_t from) { +CHAI_INLINE CHAI_HOST_DEVICE ManagedArray& ManagedArray::operator= (std::nullptr_t from) +{ m_active_pointer = from; + m_active_base_pointer = from; m_elems = 0; + m_is_slice = false; return *this; } +template +CHAI_INLINE CHAI_HOST_DEVICE bool ManagedArray::operator==( + ManagedArray& rhs) const +{ + return (m_active_pointer == rhs.m_active_pointer); +} + +template +CHAI_INLINE CHAI_HOST_DEVICE bool ManagedArray::operator!=( + ManagedArray& rhs) const +{ + return (m_active_pointer != rhs.m_active_pointer); +} + + +template +CHAI_INLINE CHAI_HOST_DEVICE bool ManagedArray::operator==(T* from) const +{ + return m_active_pointer == from; +} + +template +CHAI_INLINE CHAI_HOST_DEVICE bool ManagedArray::operator!=(T* from) const +{ + return m_active_pointer != from; +} + +template +CHAI_INLINE CHAI_HOST_DEVICE bool ManagedArray::operator==(std::nullptr_t from) const +{ + return m_active_pointer == from; +} +template +CHAI_INLINE CHAI_HOST_DEVICE bool ManagedArray::operator!=( + std::nullptr_t from) const +{ + return m_active_pointer != from; +} + +template +CHAI_INLINE CHAI_HOST_DEVICE ManagedArray::operator bool() const +{ + return m_active_pointer != nullptr; +} + } // end of namespace chai #endif // CHAI_ManagedArray_thin_INL diff --git a/src/chai/managed_ptr.hpp b/src/chai/managed_ptr.hpp index 39dca4b8..49ffa0e8 100644 --- a/src/chai/managed_ptr.hpp +++ b/src/chai/managed_ptr.hpp @@ -173,7 +173,7 @@ namespace chai { /// /// Default constructor. /// - constexpr managed_ptr() noexcept = default; + CHAI_HOST_DEVICE constexpr managed_ptr() noexcept = default; /// /// @author Alan Dayton From 10c8978678d4f612646823d125e8cae25baf747d Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Tue, 7 Apr 2020 11:05:23 -0700 Subject: [PATCH 24/70] set by value, add DISAMBIGUATE to construct from raw pointer constructor --- src/chai/ManagedArray.hpp | 24 +++++++++++++++++++----- src/chai/ManagedArray.inl | 4 ++-- src/chai/ManagedArray_thin.inl | 4 ++-- 3 files changed, 23 insertions(+), 9 deletions(-) diff --git a/src/chai/ManagedArray.hpp b/src/chai/ManagedArray.hpp index 72b04789..57286a7b 100644 --- a/src/chai/ManagedArray.hpp +++ b/src/chai/ManagedArray.hpp @@ -35,6 +35,18 @@ class CHAICopyable { }; +/*! + * \class CHAIDISAMBIGUATE + * + * \brief Type to disambiguate otherwise ambiguous constructors. + * + */ +class CHAIDISAMBIGUATE +{ +public: + CHAI_HOST_DEVICE CHAIDISAMBIGUATE(){}; + CHAI_HOST_DEVICE ~CHAIDISAMBIGUATE(){}; +}; /*! * \class ManagedArray * @@ -156,7 +168,7 @@ class ManagedArray : public CHAICopyable CHAI_HOST void move(ExecutionSpace space); - CHAI_HOST_DEVICE ManagedArray slice(size_t begin, size_t elems=-1) const; + CHAI_HOST_DEVICE ManagedArray slice(size_t begin, size_t elems=(size_t)-1) const; /*! * \brief Return reference to i-th element of the ManagedArray. * @@ -241,7 +253,7 @@ class ManagedArray : public CHAICopyable * \param val Source location of the value * \tparam T The type of data value in ManagedArray. */ - CHAI_HOST_DEVICE void set(size_t i, T& val) const; + CHAI_HOST_DEVICE void set(size_t i, T val) const; /*! * \brief Increment the value of element i in the ManagedArray. @@ -277,8 +289,10 @@ class ManagedArray : public CHAICopyable * \param data Raw pointer to data. * \param enable Boolean argument (unused) added to differentiate constructor. */ - template - CHAI_HOST_DEVICE ManagedArray(T* data, bool test = Q); + template + CHAI_HOST_DEVICE ManagedArray(T* data, + CHAIDISAMBIGUATE test = CHAIDISAMBIGUATE(), + bool foo = Q); #endif @@ -425,7 +439,7 @@ CHAI_INLINE CHAI_HOST_DEVICE ManagedArray ManagedArray::slice( size_t offs { ManagedArray slice; slice.m_resource_manager = m_resource_manager; - if (elems == -1) { + if (elems == (size_t) -1) { elems = size() - offset; } if (offset + elems > size()) { diff --git a/src/chai/ManagedArray.inl b/src/chai/ManagedArray.inl index c16bb56f..009ae376 100644 --- a/src/chai/ManagedArray.inl +++ b/src/chai/ManagedArray.inl @@ -253,7 +253,7 @@ typename ManagedArray::T_non_const ManagedArray::pick(size_t i) const { template CHAI_INLINE -CHAI_HOST_DEVICE void ManagedArray::set(size_t i, T& val) const { +CHAI_HOST_DEVICE void ManagedArray::set(size_t i, T val) const { #if !defined(__CUDA_ARCH__) && !defined(__HIP_DEVICE_COMPILE__) #if defined(CHAI_ENABLE_UM) if(m_pointer_record->m_pointers[UM] == m_active_pointer) { @@ -368,7 +368,7 @@ CHAI_HOST_DEVICE ManagedArray::operator T*() const { template template CHAI_INLINE -CHAI_HOST_DEVICE ManagedArray::ManagedArray(T* data, bool ) : +CHAI_HOST_DEVICE ManagedArray::ManagedArray(T* data, CHAIDISAMBIGUATE, bool ) : m_active_pointer(data), m_active_base_pointer(data), #if !defined(__CUDA_ARCH__) && !defined(__HIP_DEVICE_COMPILE__) diff --git a/src/chai/ManagedArray_thin.inl b/src/chai/ManagedArray_thin.inl index 834eef3f..a9fc1c93 100644 --- a/src/chai/ManagedArray_thin.inl +++ b/src/chai/ManagedArray_thin.inl @@ -193,7 +193,7 @@ typename ManagedArray::T_non_const ManagedArray::pick(size_t i) const { template CHAI_INLINE -CHAI_HOST_DEVICE void ManagedArray::set(size_t i, T& val) const { +CHAI_HOST_DEVICE void ManagedArray::set(size_t i, T val) const { #if !defined(__CUDA_ARCH__) && defined(CHAI_ENABLE_UM) cudaDeviceSynchronize(); #endif @@ -242,7 +242,7 @@ CHAI_HOST_DEVICE ManagedArray::operator T*() const { template template CHAI_INLINE -CHAI_HOST_DEVICE ManagedArray::ManagedArray(T* data, bool test) : +CHAI_HOST_DEVICE ManagedArray::ManagedArray(T* data, CHAIDISAMBIGUATE, bool) : m_active_pointer(data), m_resource_manager(ArrayManager::getInstance()), m_elems(m_resource_manager->getSize(m_active_pointer)), From c1ee1ac8b8e08205b32ae49922b4922e23e8fe69 Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Tue, 7 Apr 2020 11:15:35 -0700 Subject: [PATCH 25/70] reallocate(0) fixes, MA(record,space) fix for device usage. --- src/chai/ManagedArray.inl | 27 +++++++++++++++++++-------- 1 file changed, 19 insertions(+), 8 deletions(-) diff --git a/src/chai/ManagedArray.inl b/src/chai/ManagedArray.inl index 009ae376..1496e0d7 100644 --- a/src/chai/ManagedArray.inl +++ b/src/chai/ManagedArray.inl @@ -117,12 +117,15 @@ CHAI_INLINE CHAI_HOST_DEVICE ManagedArray::ManagedArray(PointerRecord* record, ExecutionSpace space): m_active_pointer(static_cast(record->m_pointers[space])), m_active_base_pointer(static_cast(record->m_pointers[space])), - m_resource_manager(ArrayManager::getInstance()), + m_resource_manager(nullptr), m_elems(record->m_size/sizeof(T)), m_offset(0), m_pointer_record(record), m_is_slice(false) { +#if !defined(__CUDA_ARCH__) + m_resource_manager = ArrayManager::getInstance(); +#endif } @@ -189,15 +192,23 @@ CHAI_INLINE CHAI_HOST void ManagedArray::reallocate(size_t elems) { if(!m_is_slice) { - CHAI_LOG(Debug, "Reallocating array of size " << m_elems << " with new size" << elems); + if (elems > 0) { + if (m_elems == 0 && m_active_base_pointer == nullptr) { + return allocate(elems, CPU); + } + CHAI_LOG(Debug, "Reallocating array of size " << m_elems << " with new size" << elems); - m_elems = elems; - m_active_base_pointer = - static_cast(m_resource_manager->reallocate(m_active_base_pointer, elems, - m_pointer_record)); - m_active_pointer = m_active_base_pointer; // Cannot be a slice + m_elems = elems; + m_active_base_pointer = + static_cast(m_resource_manager->reallocate(m_active_base_pointer, elems, + m_pointer_record)); + m_active_pointer = m_active_base_pointer; // Cannot be a slice - CHAI_LOG(Debug, "m_active_ptr reallocated at address: " << m_active_pointer); + CHAI_LOG(Debug, "m_active_ptr reallocated at address: " << m_active_pointer); + } + else { + this->free(); + } } } From 8cab8aa62ad8cffb2319295b6388c123d3ad3348 Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Tue, 7 Apr 2020 11:28:11 -0700 Subject: [PATCH 26/70] define behavior of pick and set when last_space is NONE, return nullptr in T* if m_elems is 0. --- src/chai/ManagedArray.inl | 19 ++++++++++++++++++- 1 file changed, 18 insertions(+), 1 deletion(-) diff --git a/src/chai/ManagedArray.inl b/src/chai/ManagedArray.inl index 1496e0d7..8e24619b 100644 --- a/src/chai/ManagedArray.inl +++ b/src/chai/ManagedArray.inl @@ -256,7 +256,14 @@ typename ManagedArray::T_non_const ManagedArray::pick(size_t i) const { return (T_non_const)(m_active_pointer[i]); } #endif - return m_resource_manager->pick(static_cast((void*)((char*)m_pointer_record->m_pointers[m_pointer_record->m_last_space]+sizeof(T)*m_offset)), i); + ExecutionSpace last_space = m_pointer_record->m_last_space; + if (last_space == NONE || last_space == CPU) { + return ((T*)m_pointer_record->m_pointers[CPU])[i+m_offset]; + } + else { + T * addr = (T*)m_pointer_record->m_pointers[last_space]; + return m_resource_manager->pick(addr, i+m_offset); + } #else return (T_non_const)(m_active_pointer[i]); #endif @@ -273,6 +280,12 @@ CHAI_HOST_DEVICE void ManagedArray::set(size_t i, T val) const { return; } #endif + + if (m_pointer_record->m_last_space == NONE) { + m_pointer_record->m_last_space = CPU; + } + + m_pointer_record->m_touched[m_pointer_record->m_last_space] = true; m_resource_manager->set(static_cast((void*)((char*)m_pointer_record->m_pointers[m_pointer_record->m_last_space]+sizeof(T)*m_offset)), i, val); #else m_active_pointer[i] = val; @@ -369,6 +382,10 @@ CHAI_HOST_DEVICE ManagedArray::operator T*() const { // Reset to whatever space we rode in on m_resource_manager->setExecutionSpace(prev_space); + if (m_elems == 0 && !m_is_slice) { + return nullptr; + } + return m_active_pointer; #else return m_active_pointer; From c939cf6de2a36a7156bfabb27c97ef3c1f147475 Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Tue, 7 Apr 2020 11:41:01 -0700 Subject: [PATCH 27/70] give PointerRecord a default constructor, avoid extra copy constructor during operator(). --- .gitignore | 3 +++ src/chai/ManagedArray.inl | 3 +-- src/chai/ManagedArray_thin.inl | 1 + src/chai/PointerRecord.hpp | 13 +++++++++++++ tests/integration/CMakeLists.txt | 1 + 5 files changed, 19 insertions(+), 2 deletions(-) diff --git a/.gitignore b/.gitignore index 2a95602e..b01cf244 100644 --- a/.gitignore +++ b/.gitignore @@ -1,3 +1,6 @@ build*/ install*/ .vscode/ +*.swp +*.swo +*.aps diff --git a/src/chai/ManagedArray.inl b/src/chai/ManagedArray.inl index 8e24619b..3822f84e 100644 --- a/src/chai/ManagedArray.inl +++ b/src/chai/ManagedArray.inl @@ -456,8 +456,7 @@ typename std::enable_if< !std::is_const::value , ManagedArray >::type () const { - return ManagedArray(const_cast(m_active_base_pointer), - m_resource_manager, m_elems, m_pointer_record); + return *reinterpret_cast *>(const_cast *>(this)); } template diff --git a/src/chai/ManagedArray_thin.inl b/src/chai/ManagedArray_thin.inl index a9fc1c93..0bc52dac 100644 --- a/src/chai/ManagedArray_thin.inl +++ b/src/chai/ManagedArray_thin.inl @@ -122,6 +122,7 @@ CHAI_HOST void ManagedArray::allocate(size_t elems, else { CHAI_LOG(Debug, "Attempted to allocate slice!"); } + m_active_base_pointer = m_active_pointer; } template diff --git a/src/chai/PointerRecord.hpp b/src/chai/PointerRecord.hpp index 9250a8e0..3308cca9 100644 --- a/src/chai/PointerRecord.hpp +++ b/src/chai/PointerRecord.hpp @@ -55,6 +55,19 @@ struct PointerRecord { UserCallback m_user_callback; int m_allocators[NUM_EXECUTION_SPACES]; + + /*! + * \brief Default constructor + * + */ + PointerRecord() : m_size(0), m_last_space(NONE) { + m_user_callback = [] (Action, ExecutionSpace, size_t) {}; + for (int space = 0; space < NUM_EXECUTION_SPACES; ++space ) { + m_pointers[space] = nullptr; + m_touched[space] = false; + m_owned[space] = true; + } + } }; } // end of namespace chai diff --git a/tests/integration/CMakeLists.txt b/tests/integration/CMakeLists.txt index d7c5c884..3dcc2c31 100644 --- a/tests/integration/CMakeLists.txt +++ b/tests/integration/CMakeLists.txt @@ -11,6 +11,7 @@ blt_list_append(TO chai_integration_test_depends ELEMENTS cuda IF ${ENABLE_CUDA} blt_list_append(TO chai_integration_test_depends ELEMENTS hip IF ${ENABLE_HIP}) blt_list_append(TO chai_integration_test_depends ELEMENTS openmp IF ${ENABLE_OPENMP}) +# ManagedArray tests blt_add_executable( NAME managed_array_tests SOURCES managed_array_tests.cpp From 946478eed1e840a854f9d22db75721091ed879e1 Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Tue, 7 Apr 2020 13:57:38 -0700 Subject: [PATCH 28/70] add some fault tolerance for when CHAI state and umpire state become inconsistent. --- src/chai/ArrayManager.cpp | 68 ++++++++++++++++++++++++++++++++++++--- src/chai/ArrayManager.hpp | 10 +++++- 2 files changed, 72 insertions(+), 6 deletions(-) diff --git a/src/chai/ArrayManager.cpp b/src/chai/ArrayManager.cpp index 1c96706c..1698fe31 100644 --- a/src/chai/ArrayManager.cpp +++ b/src/chai/ArrayManager.cpp @@ -59,26 +59,84 @@ void ArrayManager::registerPointer( std::lock_guard lock(m_mutex); auto pointer = record->m_pointers[space]; + // if we are registering a new pointer record for a pointer where there is already + // a pointer record, we assume the old record was somehow abandoned by the host + // application and trigger an ACTION_FOUND_ABANDONED callback + auto found_pointer_record_pair = m_pointer_map.find(pointer); + if (found_pointer_record_pair != m_pointer_map.end()) { + PointerRecord ** found_pointer_record_addr = found_pointer_record_pair->second; + if (found_pointer_record_addr != nullptr) { + + CHAI_LOG(Warning, "ArrayManager::registerPointer found a record for " << + pointer << " already there. Deleting abandoned pointer record."); + + PointerRecord *foundRecord = *(found_pointer_record_pair->second); + + callback(foundRecord, ACTION_FOUND_ABANDONED, space, foundRecord->m_size); + + for (int fspace = 0; fspace < NUM_EXECUTION_SPACES; ++fspace) { + foundRecord->m_pointers[fspace] = nullptr; + } + + delete foundRecord; + } + } + CHAI_LOG(Debug, "Registering " << pointer << " in space " << space); m_pointer_map.insert(pointer, record); - //record->m_last_space = space; for (int i = 0; i < NUM_EXECUTION_SPACES; i++) { if (!record->m_pointers[i]) record->m_owned[i] = true; } record->m_owned[space] = owned; -} + if (pointer) { + // if umpire already knows about this pointer, we want to make sure its records and ours + // are consistent + if (m_resource_manager.hasAllocator(pointer)) { + umpire::util::AllocationRecord *allocation_record = const_cast(m_resource_manager.findAllocationRecord(pointer)); + allocation_record->size = record->m_size; + } + // register with umpire if it's not there so that umpire can perform data migrations + else { + umpire::util::AllocationRecord new_allocation_record; + new_allocation_record.ptr = pointer; + new_allocation_record.size = record->m_size; + new_allocation_record.strategy = m_resource_manager.getAllocator(record->m_allocators[space]).getAllocationStrategy(); + + m_resource_manager.registerAllocation(pointer, new_allocation_record); + } + } +} -void ArrayManager::deregisterPointer(PointerRecord* record) +void ArrayManager::deregisterPointer(PointerRecord* record, bool deregisterFromUmpire) { std::lock_guard lock(m_mutex); for (int i = 0; i < NUM_EXECUTION_SPACES; i++) { - if (record->m_pointers[i]) m_pointer_map.erase(record->m_pointers[i]); + void * pointer = record->m_pointers[i]; + if (pointer) { + if (deregisterFromUmpire) { + m_resource_manager.deregisterAllocation(pointer); + } + m_pointer_map.erase(pointer); + } } + if (record != &s_null_record) { + delete record; + } +} - delete record; +void * ArrayManager::frontOfAllocation(void * pointer) { + if (pointer) { + if (m_resource_manager.hasAllocator(pointer)) { + auto allocation_record = m_resource_manager.findAllocationRecord(pointer); + if (allocation_record) { + return allocation_record->ptr; + } + } + } + return nullptr; } void ArrayManager::setExecutionSpace(ExecutionSpace space) diff --git a/src/chai/ArrayManager.hpp b/src/chai/ArrayManager.hpp index 1279944b..dded6cef 100644 --- a/src/chai/ArrayManager.hpp +++ b/src/chai/ArrayManager.hpp @@ -295,7 +295,15 @@ class ArrayManager /*! * \brief Deregister a PointerRecord from the ArrayManager. */ - void deregisterPointer(PointerRecord* record); + void deregisterPointer(PointerRecord* record, bool deregisterFromUmpire=false); + + /*! + * \brief Returns the front of the allocation associated with this pointer, nullptr if allocation not found. + * + * \param pointer Pointer to address of that we want the front of the allocation for. + */ + CHAISHAREDDLL_API void * frontOfAllocation(void * pointer); + /*! * \brief set the allocator for an execution space. From c3c64f9910d1e0200d13ca13c499df20e97bd6a3 Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Tue, 7 Apr 2020 14:00:37 -0700 Subject: [PATCH 29/70] add ACTION_FOUND_ABANDONED callback type. --- src/chai/Types.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/chai/Types.hpp b/src/chai/Types.hpp index 48a968ad..194c49e7 100644 --- a/src/chai/Types.hpp +++ b/src/chai/Types.hpp @@ -24,7 +24,7 @@ namespace chai typedef unsigned int uint; -enum Action { ACTION_ALLOC, ACTION_FREE, ACTION_MOVE }; +enum Action { ACTION_ALLOC, ACTION_FREE, ACTION_MOVE, ACTION_FOUND_ABANDONED }; using UserCallback = std::function; From 12498bda92a5b527a2cc48769176b5724ef6ec86 Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Tue, 7 Apr 2020 14:28:18 -0700 Subject: [PATCH 30/70] fix uninitialized m_allocators. --- src/chai/PointerRecord.hpp | 1 + 1 file changed, 1 insertion(+) diff --git a/src/chai/PointerRecord.hpp b/src/chai/PointerRecord.hpp index 3308cca9..2e96e3a7 100644 --- a/src/chai/PointerRecord.hpp +++ b/src/chai/PointerRecord.hpp @@ -66,6 +66,7 @@ struct PointerRecord { m_pointers[space] = nullptr; m_touched[space] = false; m_owned[space] = true; + m_allocators[space] = 0; } } }; From 41a0ce7761114daabe821c2c33edec71dc737817 Mon Sep 17 00:00:00 2001 From: Alan Dayton Date: Wed, 8 Apr 2020 10:40:47 -0700 Subject: [PATCH 31/70] Use ArrayManager::callback --- src/chai/ArrayManager.inl | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/src/chai/ArrayManager.inl b/src/chai/ArrayManager.inl index ab2e4475..4e930bb1 100644 --- a/src/chai/ArrayManager.inl +++ b/src/chai/ArrayManager.inl @@ -49,12 +49,12 @@ void* ArrayManager::reallocate(void* pointer, size_t elems, PointerRecord* point void* old_ptr = pointer_record->m_pointers[space]; if (old_ptr) { - pointer_record->m_user_callback(ACTION_ALLOC, ExecutionSpace(space), sizeof(T) * elems); + callback(pointer_record, ACTION_ALLOC, ExecutionSpace(space), sizeof(T) * elems); void* new_ptr = m_allocators[space]->allocate(sizeof(T)*elems); m_resource_manager.copy(new_ptr, old_ptr, num_bytes_to_copy); - pointer_record->m_user_callback(ACTION_FREE, ExecutionSpace(space), sizeof(T) * elems); + callback(pointer_record, ACTION_FREE, ExecutionSpace(space), sizeof(T) * elems); m_allocators[space]->deallocate(old_ptr); pointer_record->m_pointers[space] = new_ptr; From 7006717ba9b1980a042cc2313da1ad4083126bd5 Mon Sep 17 00:00:00 2001 From: Alan Dayton Date: Wed, 8 Apr 2020 16:01:38 -0700 Subject: [PATCH 32/70] Give more info to callbacks --- src/chai/ArrayManager.cpp | 26 +++++++++++------------ src/chai/ArrayManager.hpp | 7 +++--- src/chai/ArrayManager.inl | 24 ++++++++++++++------- src/chai/ManagedArray.hpp | 2 +- src/chai/ManagedArray.inl | 2 +- src/chai/PointerRecord.hpp | 2 +- src/chai/Types.hpp | 12 +++++++---- tests/integration/managed_array_tests.cpp | 12 +++++++---- tests/unit/array_manager_unit_tests.cpp | 6 +++--- 9 files changed, 54 insertions(+), 39 deletions(-) diff --git a/src/chai/ArrayManager.cpp b/src/chai/ArrayManager.cpp index 1698fe31..0db1d06b 100644 --- a/src/chai/ArrayManager.cpp +++ b/src/chai/ArrayManager.cpp @@ -71,8 +71,7 @@ void ArrayManager::registerPointer( pointer << " already there. Deleting abandoned pointer record."); PointerRecord *foundRecord = *(found_pointer_record_pair->second); - - callback(foundRecord, ACTION_FOUND_ABANDONED, space, foundRecord->m_size); + callback(foundRecord, ACTION_FOUND_ABANDONED, space); for (int fspace = 0; fspace < NUM_EXECUTION_SPACES; ++fspace) { foundRecord->m_pointers[fspace] = nullptr; @@ -224,9 +223,12 @@ void ArrayManager::move(PointerRecord* record, ExecutionSpace space) if (!record->m_touched[record->m_last_space]) { return; } else { - callback(record, ACTION_MOVE, space, record->m_size); - std::lock_guard lock(m_mutex); - m_resource_manager.copy(dst_pointer, src_pointer); + { + std::lock_guard lock(m_mutex); + m_resource_manager.copy(dst_pointer, src_pointer); + } + + callback(record, ACTION_MOVE, space); } resetTouch(record); @@ -239,8 +241,8 @@ void ArrayManager::allocate( auto size = pointer_record->m_size; auto alloc = m_resource_manager.getAllocator(pointer_record->m_allocators[space]); - callback(pointer_record, ACTION_ALLOC, space, size); - pointer_record->m_pointers[space] = alloc.allocate(size); + pointer_record->m_pointers[space] = alloc.allocate(size); + callback(pointer_record, ACTION_ALLOC, space); registerPointer(pointer_record, space); @@ -260,8 +262,7 @@ void ArrayManager::free(PointerRecord* pointer_record, ExecutionSpace spaceToFre if (space_ptr == pointer_record->m_pointers[UM]) { callback(pointer_record, ACTION_FREE, - ExecutionSpace(UM), - pointer_record->m_size); + ExecutionSpace(UM)); { std::lock_guard lock(m_mutex); m_pointer_map.erase(space_ptr); @@ -279,8 +280,7 @@ void ArrayManager::free(PointerRecord* pointer_record, ExecutionSpace spaceToFre #endif callback(pointer_record, ACTION_FREE, - ExecutionSpace(space), - pointer_record->m_size); + ExecutionSpace(space)); { std::lock_guard lock(m_mutex); m_pointer_map.erase(space_ptr); @@ -354,7 +354,7 @@ PointerRecord* ArrayManager::makeManaged(void* pointer, pointer_record->m_pointers[space] = pointer; pointer_record->m_owned[space] = owned; pointer_record->m_size = size; - pointer_record->m_user_callback = [](Action, ExecutionSpace, size_t) {}; + pointer_record->m_user_callback = [] (const PointerRecord*, Action, ExecutionSpace) {}; registerPointer(pointer_record, space, owned); @@ -373,7 +373,7 @@ PointerRecord* ArrayManager::deepCopyRecord(PointerRecord const* record) PointerRecord* copy = new PointerRecord{}; const size_t size = record->m_size; copy->m_size = size; - copy->m_user_callback = [](Action, ExecutionSpace, size_t) {}; + copy->m_user_callback = [] (const PointerRecord*, Action, ExecutionSpace) {}; const ExecutionSpace last_space = record->m_last_space; diff --git a/src/chai/ArrayManager.hpp b/src/chai/ArrayManager.hpp index dded6cef..37aea25b 100644 --- a/src/chai/ArrayManager.hpp +++ b/src/chai/ArrayManager.hpp @@ -329,12 +329,11 @@ class ArrayManager * \param space The space in which the event occurred * \param size The number of bytes in the array associated with this pointer record */ - inline void callback(PointerRecord* record, + inline void callback(const PointerRecord* record, Action action, - ExecutionSpace space, - size_t size) const { + ExecutionSpace space) const { if (m_callbacks_active && record) { - record->m_user_callback(action, space, size); + record->m_user_callback(record, action, space); } } diff --git a/src/chai/ArrayManager.inl b/src/chai/ArrayManager.inl index 4e930bb1..d7b0b214 100644 --- a/src/chai/ArrayManager.inl +++ b/src/chai/ArrayManager.inl @@ -36,35 +36,43 @@ void* ArrayManager::reallocate(void* pointer, size_t elems, PointerRecord* point } for (int space = CPU; space < NUM_EXECUTION_SPACES; ++space) { - if(!pointer_record->m_owned[space]) { + if (!pointer_record->m_owned[space]) { CHAI_LOG(Debug, "Cannot reallocate unowned pointer"); return pointer_record->m_pointers[my_space]; } } + // Call callback with ACTION_FREE before changing the size + for (int space = CPU; space < NUM_EXECUTION_SPACES; ++space) { + if (pointer_record->m_pointers[space]) { + callback(pointer_record, ACTION_FREE, ExecutionSpace(space)); + } + } + + // Update the pointer record size + size_t old_size = pointer_record->m_size; + size_t new_size = sizeof(T) * elems; + pointer_record->m_size = new_size; + // only copy however many bytes overlap - size_t num_bytes_to_copy = std::min(sizeof(T)*elems, pointer_record->m_size); + size_t num_bytes_to_copy = std::min(old_size, new_size); for (int space = CPU; space < NUM_EXECUTION_SPACES; ++space) { void* old_ptr = pointer_record->m_pointers[space]; if (old_ptr) { - callback(pointer_record, ACTION_ALLOC, ExecutionSpace(space), sizeof(T) * elems); - void* new_ptr = m_allocators[space]->allocate(sizeof(T)*elems); - + void* new_ptr = m_allocators[space]->allocate(new_size); m_resource_manager.copy(new_ptr, old_ptr, num_bytes_to_copy); - - callback(pointer_record, ACTION_FREE, ExecutionSpace(space), sizeof(T) * elems); m_allocators[space]->deallocate(old_ptr); pointer_record->m_pointers[space] = new_ptr; + callback(pointer_record, ACTION_ALLOC, ExecutionSpace(space)); m_pointer_map.erase(old_ptr); m_pointer_map.insert(new_ptr, pointer_record); } } - pointer_record->m_size = sizeof(T) * elems; return pointer_record->m_pointers[my_space]; } diff --git a/src/chai/ManagedArray.hpp b/src/chai/ManagedArray.hpp index 57286a7b..4eb6cd55 100644 --- a/src/chai/ManagedArray.hpp +++ b/src/chai/ManagedArray.hpp @@ -128,7 +128,7 @@ class ManagedArray : public CHAICopyable CHAI_HOST void allocate(size_t elems, ExecutionSpace space = CPU, UserCallback const& cback = - [](Action, ExecutionSpace, size_t) {}); + [] (const PointerRecord*, Action, ExecutionSpace) {}); /*! * \brief Reallocate data for the ManagedArray. diff --git a/src/chai/ManagedArray.inl b/src/chai/ManagedArray.inl index 3822f84e..63a36abc 100644 --- a/src/chai/ManagedArray.inl +++ b/src/chai/ManagedArray.inl @@ -28,7 +28,7 @@ CHAI_HOST_DEVICE ManagedArray::ManagedArray(): m_pointer_record = new PointerRecord{}; m_pointer_record->m_size = 0; - m_pointer_record->m_user_callback = [](Action, ExecutionSpace, size_t) {}; + m_pointer_record->m_user_callback = [] (const PointerRecord*, Action, ExecutionSpace) {}; for (int space = CPU; space < NUM_EXECUTION_SPACES; space++) { m_pointer_record->m_allocators[space] = diff --git a/src/chai/PointerRecord.hpp b/src/chai/PointerRecord.hpp index 2e96e3a7..49c0ce5e 100644 --- a/src/chai/PointerRecord.hpp +++ b/src/chai/PointerRecord.hpp @@ -61,7 +61,7 @@ struct PointerRecord { * */ PointerRecord() : m_size(0), m_last_space(NONE) { - m_user_callback = [] (Action, ExecutionSpace, size_t) {}; + m_user_callback = [] (const PointerRecord*, Action, ExecutionSpace) {}; for (int space = 0; space < NUM_EXECUTION_SPACES; ++space ) { m_pointers[space] = nullptr; m_touched[space] = false; diff --git a/src/chai/Types.hpp b/src/chai/Types.hpp index 194c49e7..eb087777 100644 --- a/src/chai/Types.hpp +++ b/src/chai/Types.hpp @@ -7,8 +7,12 @@ #ifndef CHAI_Types_HPP #define CHAI_Types_HPP +// Std library headers #include +// CHAI headers +#include "chai/ExecutionSpaces.hpp" + #if defined(_WIN32) && !defined(CHAISTATICLIB) #ifdef CHAISHAREDDLL_EXPORTS #define CHAISHAREDDLL_API __declspec(dllexport) @@ -21,13 +25,13 @@ namespace chai { + struct PointerRecord; -typedef unsigned int uint; - -enum Action { ACTION_ALLOC, ACTION_FREE, ACTION_MOVE, ACTION_FOUND_ABANDONED }; + typedef unsigned int uint; -using UserCallback = std::function; + enum Action { ACTION_ALLOC, ACTION_FREE, ACTION_MOVE, ACTION_FOUND_ABANDONED }; + using UserCallback = std::function; } // end of namespace chai diff --git a/tests/integration/managed_array_tests.cpp b/tests/integration/managed_array_tests.cpp index 95d40bea..66c01fe2 100644 --- a/tests/integration/managed_array_tests.cpp +++ b/tests/integration/managed_array_tests.cpp @@ -798,7 +798,8 @@ GPU_TEST(ManagedArray, UserCallback) chai::ManagedArray array; array.allocate(20, chai::CPU, - [&](chai::Action act, chai::ExecutionSpace s, size_t bytes) { + [&] (const chai::PointerRecord* record, chai::Action act, chai::ExecutionSpace s) { + const size_t bytes = record->m_size; printf("cback: act=%d, space=%d, bytes=%ld\n", (int)act, (int)s, (long)bytes); if (act == chai::ACTION_MOVE) { @@ -842,8 +843,9 @@ GPU_TEST(ManagedArray, CallBackConst) int num_h2d = 0; int num_d2h = 0; - auto callBack = [&](chai::Action act, chai::ExecutionSpace s, size_t bytes) + auto callBack = [&](const chai::PointerRecord* record, chai::Action act, chai::ExecutionSpace s) { + const size_t bytes = record->m_size; printf("cback: act=%d, space=%d, bytes=%ld\n", (int) act, (int) s, (long) bytes); if (act == chai::ACTION_MOVE) { @@ -898,8 +900,9 @@ GPU_TEST(ManagedArray, CallBackConstArray) int num_h2d = 0; int num_d2h = 0; - auto callBack = [&](chai::Action act, chai::ExecutionSpace s, size_t bytes) + auto callBack = [&] (const chai::PointerRecord* record, chai::Action act, chai::ExecutionSpace s) { + const size_t bytes = record->m_size; printf("cback: act=%d, space=%d, bytes=%ld\n", (int) act, (int) s, (long) bytes); if (act == chai::ACTION_MOVE) { @@ -997,8 +1000,9 @@ GPU_TEST(ManagedArray, CallBackConstArrayConst) int num_h2d = 0; int num_d2h = 0; - auto callBack = [&](chai::Action act, chai::ExecutionSpace s, size_t bytes) + auto callBack = [&] (const chai::PointerRecord* record, chai::Action act, chai::ExecutionSpace s) { + const size_t bytes = record->m_size; printf("cback: act=%d, space=%d, bytes=%ld\n", (int) act, (int) s, (long) bytes); if (act == chai::ACTION_MOVE) { diff --git a/tests/unit/array_manager_unit_tests.cpp b/tests/unit/array_manager_unit_tests.cpp index b50d0cce..e86fac5f 100644 --- a/tests/unit/array_manager_unit_tests.cpp +++ b/tests/unit/array_manager_unit_tests.cpp @@ -88,7 +88,7 @@ TEST(ArrayManager, controlCallbacks) // Allocate one array and set a callback size_t sizeOfArray = 5; chai::ManagedArray array1(sizeOfArray, chai::CPU); - array1.setUserCallback([&] (chai::Action, chai::ExecutionSpace, std::size_t) { + array1.setUserCallback([&] (const chai::PointerRecord*, chai::Action, chai::ExecutionSpace) { callbacksAreOn = true; }); @@ -104,7 +104,7 @@ TEST(ArrayManager, controlCallbacks) // Allocate another array and set a callback chai::ManagedArray array2(sizeOfArray, chai::CPU); - array2.setUserCallback([&] (chai::Action, chai::ExecutionSpace, std::size_t) { + array2.setUserCallback([&] (const chai::PointerRecord*, chai::Action, chai::ExecutionSpace) { callbacksAreOn = true; }); @@ -120,7 +120,7 @@ TEST(ArrayManager, controlCallbacks) // Allocate a third array and set a callback chai::ManagedArray array3(sizeOfArray, chai::CPU); - array3.setUserCallback([&] (chai::Action, chai::ExecutionSpace, std::size_t) { + array3.setUserCallback([&] (const chai::PointerRecord*, chai::Action, chai::ExecutionSpace) { callbacksAreOn = true; }); From 3ae0b417e4bfbcd0d4232dda703b5006e9cbfe09 Mon Sep 17 00:00:00 2001 From: Alan Dayton Date: Thu, 9 Apr 2020 08:58:20 -0700 Subject: [PATCH 33/70] Add a global user callback --- src/chai/ArrayManager.cpp | 5 +++++ src/chai/ArrayManager.hpp | 34 ++++++++++++++++++++++++++++++---- 2 files changed, 35 insertions(+), 4 deletions(-) diff --git a/src/chai/ArrayManager.cpp b/src/chai/ArrayManager.cpp index 0db1d06b..a0ec5f23 100644 --- a/src/chai/ArrayManager.cpp +++ b/src/chai/ArrayManager.cpp @@ -333,6 +333,11 @@ void ArrayManager::setUserCallback(void* pointer, UserCallback const& f) pointer_record->m_user_callback = f; } +void ArrayManager::setUserCallback(UserCallback const& f) +{ + m_user_callback = f; +} + PointerRecord* ArrayManager::getPointerRecord(void* pointer) { std::lock_guard lock(m_mutex); diff --git a/src/chai/ArrayManager.hpp b/src/chai/ArrayManager.hpp index 37aea25b..d0e96f8f 100644 --- a/src/chai/ArrayManager.hpp +++ b/src/chai/ArrayManager.hpp @@ -176,10 +176,17 @@ class ArrayManager bool owned); /*! - * \brief Assign a user-defined callback triggerd upon memory operations. + * \brief Assign a user-defined callback triggered upon memory operations. + * This callback applies to a single ManagedArray. */ CHAISHAREDDLL_API void setUserCallback(void* pointer, UserCallback const& f); + /*! + * \brief Assign a user-defined callback triggered upon memory operations. + * This callback applies to all ManagedArrays. + */ + CHAISHAREDDLL_API void setUserCallback(UserCallback const& f); + /*! * \brief Set touched to false in all spaces for the given PointerRecord. * @@ -332,8 +339,16 @@ class ArrayManager inline void callback(const PointerRecord* record, Action action, ExecutionSpace space) const { - if (m_callbacks_active && record) { - record->m_user_callback(record, action, space); + if (m_callbacks_active) { + // Callback for this ManagedArray only + if (record && record->m_user_callback) { + record->m_user_callback(record, action, space); + } + + // Callback for all ManagedArrays + if (m_user_callback) { + m_user_callback(record, action, space); + } } } @@ -343,7 +358,7 @@ class ArrayManager ExecutionSpace m_current_execution_space; /** - * Default space for new allocations + * Default space for new allocations. */ ExecutionSpace m_default_allocation_space; @@ -358,10 +373,21 @@ class ArrayManager */ umpire::Allocator* m_allocators[NUM_EXECUTION_SPACES]; + /*! + * \brief The umpire resource manager. + */ umpire::ResourceManager& m_resource_manager; + /*! + * \brief Used for thread-safe operations. + */ mutable std::mutex m_mutex; + /*! + * \brief A callback triggered upon memory operations on all ManagedArrays. + */ + UserCallback m_user_callback; + /*! * \brief Controls whether or not callbacks are called. */ From 7da779941132195eccaa5b36ec76b62e751515c5 Mon Sep 17 00:00:00 2001 From: Alan Dayton Date: Thu, 9 Apr 2020 09:13:54 -0700 Subject: [PATCH 34/70] Add a test for the global callback --- tests/unit/array_manager_unit_tests.cpp | 48 +++++++++++++++++++++++++ 1 file changed, 48 insertions(+) diff --git a/tests/unit/array_manager_unit_tests.cpp b/tests/unit/array_manager_unit_tests.cpp index e86fac5f..5ee77714 100644 --- a/tests/unit/array_manager_unit_tests.cpp +++ b/tests/unit/array_manager_unit_tests.cpp @@ -132,4 +132,52 @@ TEST(ArrayManager, controlCallbacks) ASSERT_TRUE(callbacksAreOn); } +/*! + * \brief Tests to see if global callback can be turned on or off + */ +TEST(ArrayManager, controlGlobalCallback) +{ + // First check that callbacks are turned on by default + chai::ArrayManager* arrayManager = chai::ArrayManager::getInstance(); + + // Variable for testing if callbacks are on or off + bool callbacksAreOn = false; + + // Set a global callback + arrayManager->setUserCallback([&] (const chai::PointerRecord*, chai::Action, chai::ExecutionSpace) { + callbacksAreOn = true; + }); + + // Allocate an array and make sure the callback was called + size_t sizeOfArray = 5; + chai::ManagedArray array(sizeOfArray, chai::CPU); + ASSERT_TRUE(callbacksAreOn); + + // Now turn off callbacks + arrayManager->disableCallbacks(); + + // Reset the variable for testing if callbacks are on or off + callbacksAreOn = false; + + // Realloc the array and make sure the callback was NOT called + array.reallocate(2 * sizeOfArray); + ASSERT_FALSE(callbacksAreOn); + + // Now make sure the order doesn't matter for when the callback is set compared + // to when callbacks are enabled + arrayManager->setUserCallback([&] (const chai::PointerRecord*, chai::Action, chai::ExecutionSpace) { + callbacksAreOn = true; + }); + + // Reset the variable for testing if callbacks are on or off + callbacksAreOn = false; + + // Turn on callbacks + arrayManager->enableCallbacks(); + + // Make sure the callback is called + array.free(); + ASSERT_TRUE(callbacksAreOn); +} + #endif // !CHAI_DISABLE_RM From 6d5a2db56b21c8839ddbbd491a74288547908c0d Mon Sep 17 00:00:00 2001 From: Alan Dayton Date: Thu, 9 Apr 2020 09:26:13 -0700 Subject: [PATCH 35/70] Add a test for callbacks --- tests/integration/managed_array_tests.cpp | 14 ++++++++++++++ 1 file changed, 14 insertions(+) diff --git a/tests/integration/managed_array_tests.cpp b/tests/integration/managed_array_tests.cpp index 66c01fe2..f6496682 100644 --- a/tests/integration/managed_array_tests.cpp +++ b/tests/integration/managed_array_tests.cpp @@ -53,6 +53,20 @@ TEST(ManagedArray, Const) array.free(); } + +TEST(ManagedArray, UserCallbackHost) +{ + bool callbackCalled = false; + + chai::ManagedArray array(10); + array.setUserCallback([&] (const chai::PointerRecord*, chai::Action, chai::ExecutionSpace) { + callbackCalled = true; + }); + + array.free(); + ASSERT_TRUE(callbackCalled); +} + #endif TEST(ManagedArray, Slice) { From 8b2dbbe8299a9f8c39ced3d008db318382880af9 Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Thu, 9 Apr 2020 11:37:08 -0700 Subject: [PATCH 36/70] bring over remaining changes - refactors migrateInner, m_pointer_record default initializaion, adds ability to free in a particular execution space. --- .gitignore | 3 + src/chai/ArrayManager.cpp | 46 +++-- src/chai/ArrayManager.hpp | 89 +++++---- src/chai/ManagedArray.hpp | 83 ++++++-- src/chai/ManagedArray.inl | 232 ++++++++++++++-------- src/chai/ManagedArray_thin.inl | 150 +++++++------- src/chai/RajaExecutionSpacePlugin.cpp | 3 + src/chai/Types.hpp | 2 +- src/chai/config.hpp.in | 1 + tests/integration/managed_array_tests.cpp | 23 +++ 10 files changed, 403 insertions(+), 229 deletions(-) diff --git a/.gitignore b/.gitignore index b01cf244..7987e0f1 100644 --- a/.gitignore +++ b/.gitignore @@ -3,4 +3,7 @@ install*/ .vscode/ *.swp *.swo +Win32/chai.vcxproj.user +Win32/.vs/ +Win32/x64/ *.aps diff --git a/src/chai/ArrayManager.cpp b/src/chai/ArrayManager.cpp index 1698fe31..8721a92d 100644 --- a/src/chai/ArrayManager.cpp +++ b/src/chai/ArrayManager.cpp @@ -172,27 +172,34 @@ ExecutionSpace ArrayManager::getExecutionSpace() void ArrayManager::registerTouch(PointerRecord* pointer_record) { - registerTouch(pointer_record, m_current_execution_space); + if (m_current_execution_space == NONE) return; + if (pointer_record) { + registerTouch(pointer_record, m_current_execution_space); + } } void ArrayManager::registerTouch(PointerRecord* pointer_record, ExecutionSpace space) { - CHAI_LOG(Debug, pointer_record->m_pointers[space] << " touched in space " << space); + if (pointer_record) { + CHAI_LOG(Debug, pointer_record->m_pointers[space] << " touched in space " << space); - if (space != NONE) { - std::lock_guard lock(m_mutex); - pointer_record->m_touched[space] = true; - pointer_record->m_last_space = space; + if (space != NONE) { + std::lock_guard lock(m_mutex); + pointer_record->m_touched[space] = true; + pointer_record->m_last_space = space; + } } } void ArrayManager::resetTouch(PointerRecord* pointer_record) { - std::lock_guard lock(m_mutex); - for (int space = CPU; space < NUM_EXECUTION_SPACES; ++space) { - pointer_record->m_touched[space] = false; + if (pointer_record && pointer_record!= &s_null_record) { + std::lock_guard lock(m_mutex); + for (int space = CPU; space < NUM_EXECUTION_SPACES; ++space) { + pointer_record->m_touched[space] = false; + } } } @@ -208,11 +215,12 @@ void ArrayManager::move(PointerRecord* record, ExecutionSpace space) } #endif + callback(record, ACTION_CAPTURED, space, record->m_size); + if (space == record->m_last_space) { return; } - void* src_pointer = record->m_pointers[record->m_last_space]; void* dst_pointer = record->m_pointers[space]; @@ -297,7 +305,7 @@ void ArrayManager::free(PointerRecord* pointer_record, ExecutionSpace spaceToFre } else { - m_resource_manager.deregisterAllocation(pointer_record->m_pointers[space]); + m_resource_manager.deregisterAllocation(pointer_record->m_pointers[space]); } } } @@ -308,6 +316,7 @@ void ArrayManager::free(PointerRecord* pointer_record, ExecutionSpace spaceToFre } } + size_t ArrayManager::getSize(void* ptr) { // TODO @@ -349,14 +358,22 @@ PointerRecord* ArrayManager::makeManaged(void* pointer, pointer, {pointer, size, m_allocators[space]->getAllocationStrategy()}); - auto pointer_record = new PointerRecord{}; - + auto pointer_record = getPointerRecord(pointer); + if (pointer_record == &s_null_record) { + pointer_record = new PointerRecord(); + } pointer_record->m_pointers[space] = pointer; pointer_record->m_owned[space] = owned; pointer_record->m_size = size; pointer_record->m_user_callback = [](Action, ExecutionSpace, size_t) {}; + + for (int space = CPU; space < NUM_EXECUTION_SPACES; ++space) { + pointer_record->m_allocators[space] = getAllocatorId(ExecutionSpace(space)); + } - registerPointer(pointer_record, space, owned); + if (pointer && size > 0) { + registerPointer(pointer_record, space, owned); + } // TODO Is this a problem? // for (int i = 0; i < NUM_EXECUTION_SPACES; i++) { @@ -428,7 +445,6 @@ int ArrayManager::getAllocatorId(ExecutionSpace space) const { return m_allocators[space]->getId(); - } void ArrayManager::evict(ExecutionSpace space, ExecutionSpace destinationSpace) { diff --git a/src/chai/ArrayManager.hpp b/src/chai/ArrayManager.hpp index dded6cef..3a12a585 100644 --- a/src/chai/ArrayManager.hpp +++ b/src/chai/ArrayManager.hpp @@ -126,7 +126,9 @@ class ArrayManager * \return Pointer to the allocated memory. */ template - void* reallocate(void* pointer, size_t elems, PointerRecord* record); + void* reallocate(void* pointer, + size_t elems, + PointerRecord* record); /*! * \brief Set the default space for new ManagedArray allocations. @@ -149,17 +151,17 @@ class ArrayManager /*! * \brief Free allocation(s) associated with the given PointerRecord. - * Default (space == NONE) will free all allocations and delete + * Default (space == NONE) will free all allocations and delete * the pointer record. */ CHAISHAREDDLL_API void free(PointerRecord* pointer, ExecutionSpace space = NONE); #if defined(CHAI_ENABLE_PICK) template - T_non_const pick(T* src_ptr, size_t index); + T_non_const pick(T* src_ptr, size_t index); template - void set(T* dst_ptr, size_t index, const T& val); + void set(T* dst_ptr, size_t index, const T& val); #endif /*! @@ -171,9 +173,9 @@ class ArrayManager CHAISHAREDDLL_API size_t getSize(void* pointer); CHAISHAREDDLL_API PointerRecord* makeManaged(void* pointer, - size_t size, - ExecutionSpace space, - bool owned); + size_t size, + ExecutionSpace space, + bool owned); /*! * \brief Assign a user-defined callback triggerd upon memory operations. @@ -234,9 +236,43 @@ class ArrayManager * \brief Wraps our resource manager's copy. * */ - CHAISHAREDDLL_API void copy(void * dst, void * src, size_t size); + CHAISHAREDDLL_API void copy(void * dst, void * src, size_t size); + + /*! + * \brief Registering an allocation with the ArrayManager + * + * \param record PointerRecord of this allocation. + * \param space Space in which the pointer was allocated. + * \param owned Should the allocation be free'd by CHAI? + */ + CHAISHAREDDLL_API void registerPointer(PointerRecord* record, + ExecutionSpace space, + bool owned = true); /*! + * \brief Deregister a PointerRecord from the ArrayManager. + * + * \param record PointerRecord of allocation to deregister. + * \param deregisterFromUmpire If true, deregister from umpire as well. + */ + CHAISHAREDDLL_API void deregisterPointer(PointerRecord* record, bool deregisterFromUmpire=false); + + /*! + * \brief Returns the front of the allocation associated with this pointer, nullptr if allocation not found. + * + * \param pointer Pointer to address of that we want the front of the allocation for. + */ + CHAISHAREDDLL_API void * frontOfAllocation(void * pointer); + + /*! + * \brief set the allocator for an execution space. + * + * \param space Execution space to set the default allocator for. + * \param allocator The allocator to use for this space. Will be copied into chai. + */ + void setAllocator(ExecutionSpace space, umpire::Allocator &allocator); + + /*! * \brief Turn callbacks on. */ void enableCallbacks() { m_callbacks_active = true; } @@ -270,6 +306,7 @@ class ArrayManager */ CHAISHAREDDLL_API void evict(ExecutionSpace space, ExecutionSpace destinationSpace); + protected: /*! * \brief Construct a new ArrayManager. @@ -279,39 +316,9 @@ class ArrayManager */ ArrayManager(); -private: - - /*! - * \brief Registering an allocation with the ArrayManager - * - * \param record PointerRecord of this allocation. - * \param space Space in which the pointer was allocated. - * \param owned Should the allocation be free'd by CHAI? - */ - void registerPointer(PointerRecord* record, - ExecutionSpace space, - bool owned = true); - /*! - * \brief Deregister a PointerRecord from the ArrayManager. - */ - void deregisterPointer(PointerRecord* record, bool deregisterFromUmpire=false); - /*! - * \brief Returns the front of the allocation associated with this pointer, nullptr if allocation not found. - * - * \param pointer Pointer to address of that we want the front of the allocation for. - */ - CHAISHAREDDLL_API void * frontOfAllocation(void * pointer); - - - /*! - * \brief set the allocator for an execution space. - * - * \param space Execution space to set the default allocator for. - * \param allocator The allocator to use for this space. Will be copied into chai. - */ - void setAllocator(ExecutionSpace space, umpire::Allocator &allocator); +private: /*! * \brief Move data in PointerRecord to the corresponding ExecutionSpace. @@ -320,8 +327,8 @@ class ArrayManager * \param space */ void move(PointerRecord* record, ExecutionSpace space); - - /*! + + /*! * \brief Execute a user callback if callbacks are active * * \param record The pointer record containing the callback diff --git a/src/chai/ManagedArray.hpp b/src/chai/ManagedArray.hpp index 57286a7b..834ac0b8 100644 --- a/src/chai/ManagedArray.hpp +++ b/src/chai/ManagedArray.hpp @@ -89,9 +89,7 @@ class ManagedArray : public CHAICopyable * \param elems Number of elements in the array. * \param space Execution space in which to allocate the array. */ - CHAI_HOST_DEVICE ManagedArray( - size_t elems, - ExecutionSpace space = NONE); + CHAI_HOST_DEVICE ManagedArray(size_t elems, ExecutionSpace space = NONE); CHAI_HOST_DEVICE ManagedArray( size_t elems, @@ -130,6 +128,7 @@ class ManagedArray : public CHAICopyable UserCallback const& cback = [](Action, ExecutionSpace, size_t) {}); + /*! * \brief Reallocate data for the ManagedArray. * @@ -142,7 +141,7 @@ class ManagedArray : public CHAICopyable /*! * \brief Free all data allocated by this ManagedArray. */ - CHAI_HOST void free(); + CHAI_HOST void free(ExecutionSpace space = NONE); /*! * \brief Reset array state. @@ -166,9 +165,10 @@ class ManagedArray : public CHAICopyable */ CHAI_HOST void registerTouch(ExecutionSpace space); - CHAI_HOST void move(ExecutionSpace space); + CHAI_HOST void move(ExecutionSpace space=NONE); CHAI_HOST_DEVICE ManagedArray slice(size_t begin, size_t elems=(size_t)-1) const; + /*! * \brief Return reference to i-th element of the ManagedArray. * @@ -181,7 +181,7 @@ class ManagedArray : public CHAICopyable /*! * \brief get access to m_active_pointer - * @return a copy of m_active_pointer + * @return a copy of m_active_base_pointer */ CHAI_HOST_DEVICE T* getActiveBasePointer() const; @@ -222,6 +222,7 @@ class ManagedArray : public CHAICopyable CHAI_HOST_DEVICE ManagedArray& operator=(std::nullptr_t); + CHAI_HOST_DEVICE bool operator==(ManagedArray& rhs) const; CHAI_HOST_DEVICE bool operator!=(ManagedArray& from) const; @@ -234,6 +235,7 @@ class ManagedArray : public CHAICopyable CHAI_HOST_DEVICE explicit operator bool() const; + #if defined(CHAI_ENABLE_PICK) /*! * \brief Return the value of element i in the ManagedArray. @@ -291,8 +293,8 @@ class ManagedArray : public CHAICopyable */ template CHAI_HOST_DEVICE ManagedArray(T* data, - CHAIDISAMBIGUATE test = CHAIDISAMBIGUATE(), - bool foo = Q); + CHAIDISAMBIGUATE test = CHAIDISAMBIGUATE(), + bool foo = Q); #endif @@ -310,14 +312,13 @@ class ManagedArray : public CHAICopyable */ CHAI_HOST void setUserCallback(UserCallback const& cback) { - m_pointer_record->m_user_callback = cback; + if (m_pointer_record && m_pointer_record != &ArrayManager::s_null_record) { + m_pointer_record->m_user_callback = cback; + } } -#endif private: - CHAI_HOST void modify(size_t i, const T& val) const; - /*! * \brief Moves the inner data of a ManagedArray. * @@ -328,7 +329,7 @@ class ManagedArray : public CHAICopyable */ template ::value, typename std::enable_if::type = 0> - CHAI_HOST void moveInnerImpl(ExecutionSpace space); + CHAI_HOST void moveInnerImpl(); /*! * \brief Does nothing since the inner data type does not inherit from @@ -341,8 +342,56 @@ class ManagedArray : public CHAICopyable */ template ::value, typename std::enable_if::type = 0> - CHAI_HOST void moveInnerImpl(ExecutionSpace space); + CHAI_HOST_DEVICE void moveInnerImpl(); +#endif +public: + CHAI_HOST_DEVICE void shallowCopy(ManagedArray const& other) const + { + m_active_pointer = other.m_active_pointer; + m_active_base_pointer = other.m_active_base_pointer; + m_resource_manager = other.m_resource_manager; + m_elems = other.m_elems; + m_offset = other.m_offset; + m_pointer_record = other.m_pointer_record; + m_is_slice = other.m_is_slice; +#ifndef CHAI_DISABLE_RM +#ifndef __CUDA_ARCH__ + // if we can, ensure elems is based off the pointer_record size to protect against + // casting leading to incorrect size info in m_elems. + if (m_pointer_record != nullptr) { + m_elems = m_pointer_record->m_size / sizeof(T); + } +#endif +#endif + } + + +private: + CHAI_HOST void modify(size_t i, const T& val) const; + // The following are only used by ManagedArray.inl, but for template + // 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 + // trigger a moveInnerImpl, which expects inner values to be initialized. + template ::value, + typename std::enable_if::type = 0> + CHAI_HOST void initInner(size_t start = 0) + { + for (size_t i = start; i < m_elems; ++i) { + m_active_base_pointer[i] = nullptr; + } + } + + // Do not deep initialize if T is not a CHAICopyable. + template ::value, + typename std::enable_if::type = 0> + CHAI_HOST void initInner(size_t = 0) + { + } +#endif +protected: /*! * Currently active data pointer. */ @@ -364,9 +413,8 @@ class ManagedArray : public CHAICopyable * Pointer to PointerRecord data. */ mutable PointerRecord* m_pointer_record = nullptr; - + mutable bool m_is_slice = false; - }; /*! @@ -424,7 +472,7 @@ template ManagedArray deepCopy(ManagedArray const& array) { T* data_ptr = array.getActiveBasePointer(); - + ArrayManager* manager = ArrayManager::getInstance(); PointerRecord const* record = manager->getPointerRecord(data_ptr); @@ -465,5 +513,4 @@ CHAI_INLINE CHAI_HOST_DEVICE ManagedArray ManagedArray::slice( size_t offs #else #include "chai/ManagedArray.inl" #endif - #endif // CHAI_ManagedArray_HPP diff --git a/src/chai/ManagedArray.inl b/src/chai/ManagedArray.inl index 3822f84e..19cdc96f 100644 --- a/src/chai/ManagedArray.inl +++ b/src/chai/ManagedArray.inl @@ -25,15 +25,7 @@ CHAI_HOST_DEVICE ManagedArray::ManagedArray(): { #if !defined(__CUDA_ARCH__) && !defined(__HIP_DEVICE_COMPILE__) m_resource_manager = ArrayManager::getInstance(); - - m_pointer_record = new PointerRecord{}; - m_pointer_record->m_size = 0; - m_pointer_record->m_user_callback = [](Action, ExecutionSpace, size_t) {}; - - for (int space = CPU; space < NUM_EXECUTION_SPACES; space++) { - m_pointer_record->m_allocators[space] = - m_resource_manager->getAllocatorId(ExecutionSpace(space)); - } + m_pointer_record = &s_null_record; #endif } @@ -45,6 +37,7 @@ CHAI_HOST_DEVICE ManagedArray::ManagedArray( ManagedArray() { #if !defined(__CUDA_ARCH__) && !defined(__HIP_DEVICE_COMPILE__) + m_pointer_record = new PointerRecord(); int i = 0; for (auto& space : spaces) { m_pointer_record->m_allocators[space] = allocators.begin()[i++].getId(); @@ -61,11 +54,9 @@ CHAI_HOST_DEVICE ManagedArray::ManagedArray( ManagedArray() { #if !defined(__CUDA_ARCH__) && !defined(__HIP_DEVICE_COMPILE__) - m_elems = elems; - m_pointer_record->m_size = sizeof(T)*m_elems; this->allocate(elems, space); - + #if defined(CHAI_ENABLE_UM) if(space == UM) { m_pointer_record->m_pointers[CPU] = m_active_pointer; @@ -86,7 +77,6 @@ CHAI_HOST_DEVICE ManagedArray::ManagedArray( { #if !defined(__CUDA_ARCH__) && !defined(__HIP_DEVICE_COMPILE__) m_elems = elems; - m_pointer_record->m_size = sizeof(T)*elems; this->allocate(elems, space); @@ -110,6 +100,10 @@ CHAI_HOST_DEVICE ManagedArray::ManagedArray(std::nullptr_t) : m_pointer_record(nullptr), m_is_slice(false) { +#if !defined(__CUDA_ARCH__) && !defined(__HIP_DEVICE_COMPILE__) + m_resource_manager = ArrayManager::getInstance(); + m_pointer_record = &ArrayManager::s_null_record(); +#endif } template @@ -123,8 +117,11 @@ CHAI_HOST_DEVICE ManagedArray::ManagedArray(PointerRecord* record, ExecutionS m_pointer_record(record), m_is_slice(false) { -#if !defined(__CUDA_ARCH__) +#if !defined(__CUDA_ARCH__) && !defined(__HIP_DEVICE_COMPILE__) m_resource_manager = ArrayManager::getInstance(); + if (m_pointer_record == nullptr) { + m_pointer_record = &ArrayManager::s_null_record; + } #endif } @@ -141,10 +138,13 @@ CHAI_HOST_DEVICE ManagedArray::ManagedArray(ManagedArray const& other): m_is_slice(other.m_is_slice) { #if !defined(__CUDA_ARCH__) && !defined(__HIP_DEVICE_COMPILE__) - if (other.m_pointer_record != nullptr) { - m_elems = other.m_pointer_record->m_size/sizeof(T); + m_elems = other.m_pointer_record->m_size/sizeof(T); + if (m_active_base_pointer) { + /// Move nested ManagedArrays first, so they are working with a valid m_active_pointer for the host, + // and so the meta data associated with them are updated before we move that down. + moveInnerImpl(); + move(); } - move(m_resource_manager->getExecutionSpace()); #endif } @@ -159,8 +159,19 @@ CHAI_HOST_DEVICE ManagedArray::ManagedArray(T* data, ArrayManager* array_mana m_pointer_record(pointer_record), m_is_slice(false) { +#if !defined(__CUDA_ARCH__) && !defined(__HIP_DEVICE_COMPILE__) + if (m_resource_manager == nullptr) { + m_resource_manager = ArrayManager::getInstance(); + } + if (m_pointer_record == &chai::ArrayManager::s_null_record || m_pointer_record==nullptr) { + bool owned = true; + m_pointer_record = m_resource_manager->makeManaged((void *) data, sizeof(T)*m_elems,ExecutionSpace(CPU),true); + } + registerTouch(CPU); +#endif } + template CHAI_HOST void ManagedArray::allocate( size_t elems, @@ -168,25 +179,38 @@ CHAI_HOST void ManagedArray::allocate( const UserCallback& cback) { if(!m_is_slice) { - CHAI_LOG(Debug, "Allocating array of size " << elems << " in space " << space); + if (elems > 0) { + CHAI_LOG(Debug, "Allocating array of size " << elems << " in space " << space); - if (space == NONE) { - space = m_resource_manager->getDefaultAllocationSpace(); - } + if (space == NONE) { + space = m_resource_manager->getDefaultAllocationSpace(); + } + if (m_pointer_record == &chai::ArrayManager::s_null_record) { + m_pointer_record = m_resource_manager->makeManaged((void *) m_active_base_pointer,m_elems*sizeof(T),CPU,true); + } - setUserCallback(cback); - m_elems = elems; - m_pointer_record->m_size = sizeof(T)*elems; + m_pointer_record->m_user_callback = cback; + m_elems = elems; + m_pointer_record->m_size = sizeof(T)*elems; - m_resource_manager->allocate(m_pointer_record, space); + m_resource_manager->allocate(m_pointer_record, space); - m_active_base_pointer = static_cast(m_pointer_record->m_pointers[space]); - m_active_pointer = m_active_base_pointer; // Cannot be a slice + m_active_base_pointer = static_cast(m_pointer_record->m_pointers[space]); + m_active_pointer = m_active_base_pointer; // Cannot be a slice - CHAI_LOG(Debug, "m_active_ptr allocated at address: " << m_active_pointer); + // if T is a CHAICopyable, then it is important to initialize all the + // ManagedArrays to nullptr at allocation, since it is extremely easy to + // trigger a moveInnerImpl, which expects inner values to be initialized. + initInner(); + + CHAI_LOG(Debug, "m_active_ptr allocated at address: " << m_active_pointer); + } } } + + + template CHAI_INLINE CHAI_HOST void ManagedArray::reallocate(size_t elems) @@ -197,12 +221,21 @@ CHAI_HOST void ManagedArray::reallocate(size_t elems) return allocate(elems, CPU); } CHAI_LOG(Debug, "Reallocating array of size " << m_elems << " with new size" << elems); + if (m_pointer_record == &chai::ArrayManager::s_null_record) { + m_pointer_record = m_resource_manager->makeManaged((void *)m_active_base_pointer,m_elems*sizeof(T),CPU,true); + } + size_t old_size = m_elems; m_elems = elems; m_active_base_pointer = static_cast(m_resource_manager->reallocate(m_active_base_pointer, elems, - m_pointer_record)); + m_pointer_record)); m_active_pointer = m_active_base_pointer; // Cannot be a slice + + // if T is a CHAICopyable, then it is important to initialize all the new + // ManagedArrays to nullptr at allocation, since it is extremely easy to + // trigger a moveInnerImpl, which expects inner values to be initialized. + initInner(old_size); CHAI_LOG(Debug, "m_active_ptr reallocated at address: " << m_active_pointer); } @@ -214,11 +247,26 @@ CHAI_HOST void ManagedArray::reallocate(size_t elems) template CHAI_INLINE -CHAI_HOST void ManagedArray::free() +CHAI_HOST void ManagedArray::free(ExecutionSpace space) { - if(!m_is_slice) { - m_resource_manager->free(m_pointer_record); - m_pointer_record = nullptr; + if(!m_is_slice && *this != nullptr) { + if (m_resource_manager == nullptr) { + m_resource_manager = ArrayManager::getInstance(); + } + if (m_pointer_record == &ArrayManager::s_null_record) { + m_pointer_record = m_resource_manager->makeManaged((void *)m_active_base_pointer,m_elems*sizeof(T),space,true); + } + m_resource_manager->free(m_pointer_record, space); + m_active_pointer = nullptr; + m_active_base_pointer = nullptr; + + m_elems = 0; + m_offset = 0; + // The call to m_resource_manager::free, above, has deallocated m_pointer_record if space == NONE. + if (space == NONE) { + m_pointer_record = &s_null_record; + } + m_is_slice = false; } else { CHAI_LOG(Debug, "Cannot free a slice!"); } @@ -240,6 +288,10 @@ CHAI_HOST_DEVICE size_t ManagedArray::size() const { template CHAI_INLINE CHAI_HOST void ManagedArray::registerTouch(ExecutionSpace space) { + if (m_active_pointer && (m_pointer_record == nullptr || m_pointer_record == &ArrayManager::s_null_record)) { + CHAI_LOG(Warning,"registerTouch called on ManagedArray with nullptr pointer record."); + m_pointer_record = m_resource_manager->makeManaged((void *)m_active_base_pointer,m_elems*sizeof(T),CPU,true); + } m_resource_manager->registerTouch(m_pointer_record, space); } @@ -271,7 +323,7 @@ typename ManagedArray::T_non_const ManagedArray::pick(size_t i) const { template CHAI_INLINE -CHAI_HOST_DEVICE void ManagedArray::set(size_t i, T val) const { +CHAI_HOST_DEVICE void ManagedArray::set(size_t i, T val) const { #if !defined(__CUDA_ARCH__) && !defined(__HIP_DEVICE_COMPILE__) #if defined(CHAI_ENABLE_UM) if(m_pointer_record->m_pointers[UM] == m_active_pointer) { @@ -289,7 +341,7 @@ CHAI_HOST_DEVICE void ManagedArray::set(size_t i, T val) const { m_resource_manager->set(static_cast((void*)((char*)m_pointer_record->m_pointers[m_pointer_record->m_last_space]+sizeof(T)*m_offset)), i, val); #else m_active_pointer[i] = val; - #endif + #endif // !defined(__CUDA_ARCH__) } template @@ -333,29 +385,20 @@ CHAI_INLINE CHAI_HOST void ManagedArray::move(ExecutionSpace space) { - ExecutionSpace prev_space = m_pointer_record->m_last_space; - - /* When moving from CPU to GPU we need to move the inner arrays before the outer array. */ - if (prev_space == CPU) { - moveInnerImpl(space); + if (m_resource_manager == nullptr) { + m_resource_manager = ArrayManager::getInstance(); } - - m_active_base_pointer = static_cast(m_resource_manager->move(const_cast(m_active_base_pointer), m_pointer_record, space)); - m_active_pointer = m_active_base_pointer + m_offset; - - if (!std::is_const::value) { - CHAI_LOG(Debug, "T is non-const, registering touch of pointer" << m_active_pointer); - m_resource_manager->registerTouch(m_pointer_record, space); - } - - if (space != NONE) m_pointer_record->m_last_space = space; - - /* When moving from GPU to CPU we need to move the inner arrays after the outer array. */ -#if defined(CHAI_ENABLE_CUDA) - if (space != GPU && prev_space == GPU) { - moveInnerImpl(space); + if (m_pointer_record != &ArrayManager::s_null_record) { + CHAI_LOG(Debug, "Moving " << m_active_pointer); + m_active_base_pointer = static_cast(m_resource_manager->move((void *)m_active_base_pointer, m_pointer_record, space)); + m_active_pointer = m_active_base_pointer + m_offset; + + CHAI_LOG(Debug, "Moved to " << m_active_pointer); + if (!std::is_const::value) { + CHAI_LOG(Debug, "T is non-const, registering touch of pointer" << m_active_pointer); + m_resource_manager->registerTouch(m_pointer_record); + } } -#endif } template @@ -370,17 +413,25 @@ template CHAI_INLINE CHAI_HOST_DEVICE ManagedArray::operator T*() const { #if !defined(__CUDA_ARCH__) && !defined(__HIP_DEVICE_COMPILE__) - ExecutionSpace prev_space = m_resource_manager->getExecutionSpace(); - m_resource_manager->setExecutionSpace(CPU); - auto non_const_active_base_pointer = const_cast(static_cast(m_active_base_pointer)); - m_active_base_pointer = static_cast(m_resource_manager->move(non_const_active_base_pointer, m_pointer_record)); - m_active_pointer = m_active_base_pointer + m_offset; - - m_resource_manager->registerTouch(m_pointer_record); - - - // Reset to whatever space we rode in on - m_resource_manager->setExecutionSpace(prev_space); + if (m_active_pointer) { + if (m_resource_manager == nullptr) { + m_resource_manager = ArrayManager::getInstance(); + } + if (m_pointer_record == nullptr || m_pointer_record == &ArrayManager::s_null_record) { + CHAI_LOG(Warning, "nullptr pointer_record associated with non-nullptr active_pointer") + } + ExecutionSpace prev_space = m_resource_manager->getExecutionSpace(); + m_resource_manager->setExecutionSpace(CPU); + auto non_const_active_base_pointer = const_cast(static_cast(m_active_base_pointer)); + m_active_base_pointer = static_cast(m_resource_manager->move(non_const_active_base_pointer, m_pointer_record)); + m_active_pointer = m_active_base_pointer+m_offset; + + m_resource_manager->registerTouch(m_pointer_record); + + + // Reset to whatever space we rode in on + m_resource_manager->setExecutionSpace(prev_space); + } if (m_elems == 0 && !m_is_slice) { return nullptr; @@ -402,7 +453,7 @@ CHAI_HOST_DEVICE ManagedArray::ManagedArray(T* data, CHAIDISAMBIGUATE, bool ) #if !defined(__CUDA_ARCH__) && !defined(__HIP_DEVICE_COMPILE__) m_resource_manager(ArrayManager::getInstance()), m_elems(m_resource_manager->getSize((void *)m_active_base_pointer)/sizeof(T)), - m_pointer_record(m_resource_manager->getPointerRecord(data)), + m_pointer_record(m_resource_manager->getPointerRecord((void *)data)), #else m_resource_manager(nullptr), m_elems(0), @@ -411,6 +462,12 @@ CHAI_HOST_DEVICE ManagedArray::ManagedArray(T* data, CHAIDISAMBIGUATE, bool ) m_offset(0), m_is_slice(false) { +#if !defined(__CUDA_ARCH__) && !defined(__HIP_DEVICE_COMPILE__) + + if (m_active_pointer && (m_pointer_record == &ArrayManager::s_null_record || m_active_pointer != m_pointer_record->m_pointers[CPU])) { + CHAI_LOG(Warning,"REINTEGRATED external pointer unknown by CHAI."); + } +#endif } #endif @@ -428,10 +485,13 @@ ManagedArray::getActivePointer() const return m_active_pointer; } -template +template T* -ManagedArray::getPointer(ExecutionSpace space, bool do_move) { - if (m_elems == 0 && !m_is_slice) { +ManagedArray::getPointer(ExecutionSpace space, bool do_move) { + if (m_pointer_record == nullptr || m_pointer_record == &ArrayManager::s_null_record) { + return nullptr; + } + if (m_elems == 0 && !m_is_slice) { return nullptr; } if (do_move) { @@ -444,6 +504,7 @@ ManagedArray::getPointer(ExecutionSpace space, bool do_move) { return ((T*) m_pointer_record->m_pointers[space]) + offset; } + //template //ManagedArray::operator ManagedArray< // typename std::conditional::value, @@ -478,7 +539,11 @@ ManagedArray::operator= (std::nullptr_t) { m_active_base_pointer = nullptr; m_elems = 0; m_offset = 0; + #if !defined(__CUDA_ARCH__) && !defined(__HIP_DEVICE_COMPILE__) + m_pointer_record = ArrayManager::&s_null_record; + #else m_pointer_record = nullptr; + #endif m_is_slice = false; return *this; } @@ -545,20 +610,17 @@ template::type> CHAI_INLINE CHAI_HOST void -ManagedArray::moveInnerImpl(ExecutionSpace space) { - if (space == NONE) { - return; - } - - ExecutionSpace const prev_space = m_resource_manager->getExecutionSpace(); - m_resource_manager->setExecutionSpace(space); - - T_non_const* non_const_active_base_pointer = const_cast(m_active_base_pointer); - for (int i = 0; i < size(); ++i) { - non_const_active_base_pointer[i] = T(m_active_base_pointer[i]); +ManagedArray::moveInnerImpl() +{ + int len = m_pointer_record->m_size / sizeof(T); + T * host_ptr = (T *) m_pointer_record->m_pointers[CPU]; + for (int i = 0; i < len; ++i) + { + // trigger the copy constructor + T inner = T(host_ptr[i]); + // ensure the inner type gets the state of the result of the copy + host_ptr[i].shallowCopy(inner); } - - m_resource_manager->setExecutionSpace(prev_space); } template @@ -566,7 +628,7 @@ template::type> CHAI_INLINE CHAI_HOST void -ManagedArray::moveInnerImpl(ExecutionSpace CHAI_UNUSED_ARG(space)) +ManagedArray::moveInnerImpl() { } diff --git a/src/chai/ManagedArray_thin.inl b/src/chai/ManagedArray_thin.inl index 0bc52dac..0655a990 100644 --- a/src/chai/ManagedArray_thin.inl +++ b/src/chai/ManagedArray_thin.inl @@ -16,11 +16,10 @@ namespace chai { template -CHAI_INLINE -CHAI_HOST_DEVICE ManagedArray::ManagedArray( - std::initializer_list spaces, - std::initializer_list allocators) : - ManagedArray() +CHAI_INLINE CHAI_HOST_DEVICE ManagedArray::ManagedArray( + std::initializer_list spaces, + std::initializer_list allocators) + : ManagedArray() { if (m_pointer_record) { int i = 0; @@ -56,6 +55,7 @@ CHAI_HOST_DEVICE ManagedArray::ManagedArray(size_t elems, ExecutionSpace spac this->allocate(elems, space); } + template CHAI_INLINE CHAI_HOST_DEVICE ManagedArray::ManagedArray(std::nullptr_t) @@ -155,114 +155,127 @@ CHAI_HOST void ManagedArray::reallocate(size_t new_elems) } } -template -CHAI_INLINE -CHAI_HOST void ManagedArray::free() +template +CHAI_INLINE CHAI_HOST void ManagedArray::free(ExecutionSpace space) { - if (!m_is_slice) { - #if defined(CHAI_ENABLE_UM) - cudaFree(m_active_base_pointer); - #else - ::free((void *)m_active_base_pointer); - #endif - - m_active_base_pointer = nullptr; - m_active_pointer = nullptr; - } - else { - CHAI_LOG(Debug, "tried to free slice!"); + if (space == CPU || space == NONE) { + if (!m_is_slice) { +#if defined(CHAI_ENABLE_UM) + cudaFree(m_active_pointer); +#else + ::free(m_active_pointer); +#endif + m_active_pointer = nullptr; + m_active_base_pointer = nullptr; + } else { + CHAI_LOG(Debug, "tried to free slice!"); + } } } -template -CHAI_INLINE -CHAI_HOST void ManagedArray::reset() + +template +CHAI_INLINE CHAI_HOST void ManagedArray::reset() { } #if defined(CHAI_ENABLE_PICK) -template -CHAI_INLINE -CHAI_HOST_DEVICE -typename ManagedArray::T_non_const ManagedArray::pick(size_t i) const { +template +CHAI_INLINE CHAI_HOST_DEVICE typename ManagedArray::T_non_const ManagedArray< + T>::pick(size_t i) const +{ #if !defined(__CUDA_ARCH__) && defined(CHAI_ENABLE_UM) cudaDeviceSynchronize(); #endif - return (T_non_const) m_active_pointer[i]; + return (T_non_const)m_active_pointer[i]; } -template -CHAI_INLINE -CHAI_HOST_DEVICE void ManagedArray::set(size_t i, T val) const { +template +CHAI_INLINE CHAI_HOST_DEVICE void ManagedArray::set(size_t i, T val) const +{ #if !defined(__CUDA_ARCH__) && defined(CHAI_ENABLE_UM) cudaDeviceSynchronize(); #endif - m_active_pointer[i] = val; + m_active_pointer[i] = val; } -template -CHAI_INLINE -CHAI_HOST_DEVICE void ManagedArray::incr(size_t i) const { +template +CHAI_INLINE CHAI_HOST_DEVICE void ManagedArray::incr(size_t i) const +{ #if !defined(__CUDA_ARCH__) && defined(CHAI_ENABLE_UM) cudaDeviceSynchronize(); #endif - ++m_active_pointer[i]; + ++m_active_pointer[i]; } -template -CHAI_INLINE -CHAI_HOST_DEVICE void ManagedArray::decr(size_t i) const { +template +CHAI_INLINE CHAI_HOST_DEVICE void ManagedArray::decr(size_t i) const +{ #if !defined(__CUDA_ARCH__) && defined(CHAI_ENABLE_UM) cudaDeviceSynchronize(); #endif - --m_active_pointer[i]; + --m_active_pointer[i]; } #endif -template -CHAI_INLINE -CHAI_HOST size_t ManagedArray::size() const { +template +CHAI_INLINE CHAI_HOST_DEVICE size_t ManagedArray::size() const +{ return m_elems; } -template -template -CHAI_INLINE -CHAI_HOST_DEVICE T& ManagedArray::operator[](const Idx i) const { +template +CHAI_INLINE CHAI_HOST void ManagedArray::registerTouch(ExecutionSpace) +{ +} + +template +CHAI_INLINE CHAI_HOST void ManagedArray::move(ExecutionSpace) +{ +} + +template +template +CHAI_INLINE CHAI_HOST_DEVICE T& ManagedArray::operator[](const Idx i) const +{ return m_active_pointer[i]; } #if defined(CHAI_ENABLE_IMPLICIT_CONVERSIONS) -template -CHAI_INLINE -CHAI_HOST_DEVICE ManagedArray::operator T*() const { +template +CHAI_INLINE CHAI_HOST_DEVICE ManagedArray::operator T*() const +{ return m_active_pointer; } -template -template -CHAI_INLINE -CHAI_HOST_DEVICE ManagedArray::ManagedArray(T* data, CHAIDISAMBIGUATE, bool) : +template +template +CHAI_INLINE CHAI_HOST_DEVICE +ManagedArray::ManagedArray(T* data, CHAIDISAMBIGUATE, bool) : m_active_pointer(data), - m_resource_manager(ArrayManager::getInstance()), - m_elems(m_resource_manager->getSize(m_active_pointer)), + m_active_base_pointer(data), + m_resource_manager(nullptr), + m_elems(-1), + m_pointer_record(nullptr), m_is_slice(false) { } #endif -template -template< typename U> -ManagedArray::operator -typename std::enable_if< !std::is_const::value , - ManagedArray >::type () const +template +template +ManagedArray::operator typename std:: + enable_if::value, ManagedArray >::type() const { - return ManagedArray(const_cast(m_active_pointer), m_resource_manager, m_elems, nullptr); + return ManagedArray(const_cast(m_active_pointer), + m_resource_manager, + m_elems, + nullptr); } -template -CHAI_INLINE CHAI_HOST_DEVICE ManagedArray& ManagedArray::operator= (std::nullptr_t from) +template +CHAI_INLINE CHAI_HOST_DEVICE ManagedArray& ManagedArray::operator=(std::nullptr_t from) { m_active_pointer = from; m_active_base_pointer = from; @@ -299,13 +312,12 @@ CHAI_INLINE CHAI_HOST_DEVICE bool ManagedArray::operator!=(T* from) const } template -CHAI_INLINE CHAI_HOST_DEVICE bool ManagedArray::operator==(std::nullptr_t from) const +CHAI_INLINE CHAI_HOST_DEVICE bool ManagedArray::operator==( std::nullptr_t from) const { return m_active_pointer == from; } template -CHAI_INLINE CHAI_HOST_DEVICE bool ManagedArray::operator!=( - std::nullptr_t from) const +CHAI_INLINE CHAI_HOST_DEVICE bool ManagedArray::operator!=( std::nullptr_t from) const { return m_active_pointer != from; } @@ -316,6 +328,6 @@ CHAI_INLINE CHAI_HOST_DEVICE ManagedArray::operator bool() const return m_active_pointer != nullptr; } -} // end of namespace chai +} // end of namespace chai -#endif // CHAI_ManagedArray_thin_INL +#endif // CHAI_ManagedArray_thin_INL diff --git a/src/chai/RajaExecutionSpacePlugin.cpp b/src/chai/RajaExecutionSpacePlugin.cpp index 6ba50463..0584a9eb 100644 --- a/src/chai/RajaExecutionSpacePlugin.cpp +++ b/src/chai/RajaExecutionSpacePlugin.cpp @@ -75,6 +75,7 @@ RajaExecutionSpacePlugin::postLaunch(RAJA::util::PluginContext) } } +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) @@ -96,8 +97,10 @@ RAJA::util::PluginRegistry::Add P( "RajaExecutionSpacePlugin", "Plugin to set CHAI execution space based on RAJA execution platform"); + namespace chai { void linkRajaPlugin() {} } + diff --git a/src/chai/Types.hpp b/src/chai/Types.hpp index 194c49e7..fe8a1230 100644 --- a/src/chai/Types.hpp +++ b/src/chai/Types.hpp @@ -24,7 +24,7 @@ namespace chai typedef unsigned int uint; -enum Action { ACTION_ALLOC, ACTION_FREE, ACTION_MOVE, ACTION_FOUND_ABANDONED }; +enum Action { ACTION_ALLOC, ACTION_FREE, ACTION_MOVE, ACTION_CAPTURED, ACTION_FOUND_ABANDONED }; using UserCallback = std::function; diff --git a/src/chai/config.hpp.in b/src/chai/config.hpp.in index 5a1c53ac..40ba15d5 100644 --- a/src/chai/config.hpp.in +++ b/src/chai/config.hpp.in @@ -14,6 +14,7 @@ #cmakedefine CHAI_DISABLE_RM #cmakedefine CHAI_ENABLE_UM #cmakedefine CHAI_DEBUG +#cmakedefine CHAI_ENABLE_GPU_ERROR_CHECKING #cmakedefine CHAI_ENABLE_RAJA_PLUGIN #cmakedefine CHAI_ENABLE_GPU_SIMULATION_MODE diff --git a/tests/integration/managed_array_tests.cpp b/tests/integration/managed_array_tests.cpp index 95d40bea..6e0a4e71 100644 --- a/tests/integration/managed_array_tests.cpp +++ b/tests/integration/managed_array_tests.cpp @@ -1417,6 +1417,29 @@ GPU_TEST(ManagedArray, DeviceDeepCopy) #endif #endif // defined(CHAI_ENABLE_CUDA) || defined(CHAI_ENABLE_HIP) +#ifdef CHAI_ENABLE_CUDA + +CUDA_TEST(ManagedArray, CopyConstruct) +{ + const int expectedValue = rand(); + + chai::ManagedArray array(1, chai::CPU); + array[0] = expectedValue; + + chai::ManagedArray array2 = array; + + chai::ManagedArray results(1, chai::GPU); + + forall(cuda(), 0, 1, [=] __device__ (int i) { + results[i] = array2[i]; + }); + + results.move(chai::CPU); + ASSERT_EQ(results[0], expectedValue); +} + +#endif + TEST(ManagedArray, SizeZero) { chai::ManagedArray array; From 157aa4123d4d1c66112eed3e05224c175a90df5d Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Thu, 9 Apr 2020 11:44:33 -0700 Subject: [PATCH 37/70] build fixes. --- src/chai/ManagedArray.inl | 14 +++++++------- 1 file changed, 7 insertions(+), 7 deletions(-) diff --git a/src/chai/ManagedArray.inl b/src/chai/ManagedArray.inl index 19cdc96f..99be774e 100644 --- a/src/chai/ManagedArray.inl +++ b/src/chai/ManagedArray.inl @@ -25,7 +25,7 @@ CHAI_HOST_DEVICE ManagedArray::ManagedArray(): { #if !defined(__CUDA_ARCH__) && !defined(__HIP_DEVICE_COMPILE__) m_resource_manager = ArrayManager::getInstance(); - m_pointer_record = &s_null_record; + m_pointer_record = &ArrayManager::s_null_record; #endif } @@ -102,7 +102,7 @@ CHAI_HOST_DEVICE ManagedArray::ManagedArray(std::nullptr_t) : { #if !defined(__CUDA_ARCH__) && !defined(__HIP_DEVICE_COMPILE__) m_resource_manager = ArrayManager::getInstance(); - m_pointer_record = &ArrayManager::s_null_record(); + m_pointer_record = &ArrayManager::s_null_record; #endif } @@ -163,7 +163,7 @@ CHAI_HOST_DEVICE ManagedArray::ManagedArray(T* data, ArrayManager* array_mana if (m_resource_manager == nullptr) { m_resource_manager = ArrayManager::getInstance(); } - if (m_pointer_record == &chai::ArrayManager::s_null_record || m_pointer_record==nullptr) { + if (m_pointer_record == &ArrayManager::s_null_record || m_pointer_record==nullptr) { bool owned = true; m_pointer_record = m_resource_manager->makeManaged((void *) data, sizeof(T)*m_elems,ExecutionSpace(CPU),true); } @@ -185,7 +185,7 @@ CHAI_HOST void ManagedArray::allocate( if (space == NONE) { space = m_resource_manager->getDefaultAllocationSpace(); } - if (m_pointer_record == &chai::ArrayManager::s_null_record) { + if (m_pointer_record == &ArrayManager::s_null_record) { m_pointer_record = m_resource_manager->makeManaged((void *) m_active_base_pointer,m_elems*sizeof(T),CPU,true); } @@ -221,7 +221,7 @@ CHAI_HOST void ManagedArray::reallocate(size_t elems) return allocate(elems, CPU); } CHAI_LOG(Debug, "Reallocating array of size " << m_elems << " with new size" << elems); - if (m_pointer_record == &chai::ArrayManager::s_null_record) { + if (m_pointer_record == &ArrayManager::s_null_record) { m_pointer_record = m_resource_manager->makeManaged((void *)m_active_base_pointer,m_elems*sizeof(T),CPU,true); } size_t old_size = m_elems; @@ -264,7 +264,7 @@ CHAI_HOST void ManagedArray::free(ExecutionSpace space) m_offset = 0; // The call to m_resource_manager::free, above, has deallocated m_pointer_record if space == NONE. if (space == NONE) { - m_pointer_record = &s_null_record; + m_pointer_record = &ArrayManager::s_null_record; } m_is_slice = false; } else { @@ -540,7 +540,7 @@ ManagedArray::operator= (std::nullptr_t) { m_elems = 0; m_offset = 0; #if !defined(__CUDA_ARCH__) && !defined(__HIP_DEVICE_COMPILE__) - m_pointer_record = ArrayManager::&s_null_record; + m_pointer_record = &ArrayManager::s_null_record; #else m_pointer_record = nullptr; #endif From 0929a7eaf58bf788aa3ff2a19c3778305da86346 Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Thu, 9 Apr 2020 11:56:57 -0700 Subject: [PATCH 38/70] fix build error from merge. --- src/chai/ArrayManager.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/chai/ArrayManager.cpp b/src/chai/ArrayManager.cpp index 1c9ba007..00fada3c 100644 --- a/src/chai/ArrayManager.cpp +++ b/src/chai/ArrayManager.cpp @@ -214,7 +214,7 @@ void ArrayManager::move(PointerRecord* record, ExecutionSpace space) } #endif - callback(record, ACTION_CAPTURED, space, record->m_size); + callback(record, ACTION_CAPTURED, space); if (space == record->m_last_space) { return; From 73ff6f4924df50a4c810e30e6353b42542787165 Mon Sep 17 00:00:00 2001 From: robinson96 Date: Thu, 9 Apr 2020 13:36:15 -0700 Subject: [PATCH 39/70] Update tests/integration/managed_array_tests.cpp CUDA_TEST should be GPU_TEST Co-Authored-By: David Beckingsale --- tests/integration/managed_array_tests.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/integration/managed_array_tests.cpp b/tests/integration/managed_array_tests.cpp index c2612d9d..415026ff 100644 --- a/tests/integration/managed_array_tests.cpp +++ b/tests/integration/managed_array_tests.cpp @@ -1423,7 +1423,7 @@ GPU_TEST(ManagedArray, DeviceDeepCopy) #ifdef CHAI_ENABLE_CUDA -CUDA_TEST(ManagedArray, CopyConstruct) +GPU_TEST(ManagedArray, CopyConstruct) { const int expectedValue = rand(); From 1812b3c07f5b0a17e8510be057e64fa4cf30b90b Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Thu, 9 Apr 2020 14:03:05 -0700 Subject: [PATCH 40/70] cuda->gpu --- tests/integration/managed_array_tests.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/integration/managed_array_tests.cpp b/tests/integration/managed_array_tests.cpp index 415026ff..a4a4e457 100644 --- a/tests/integration/managed_array_tests.cpp +++ b/tests/integration/managed_array_tests.cpp @@ -1434,7 +1434,7 @@ GPU_TEST(ManagedArray, CopyConstruct) chai::ManagedArray results(1, chai::GPU); - forall(cuda(), 0, 1, [=] __device__ (int i) { + forall(gpu(), 0, 1, [=] __device__ (int i) { results[i] = array2[i]; }); From daefac147efe47eee590ec9008611990045c16b6 Mon Sep 17 00:00:00 2001 From: Alan Dayton Date: Thu, 9 Apr 2020 14:40:28 -0700 Subject: [PATCH 41/70] Rename global callback setter --- src/chai/ArrayManager.cpp | 2 +- src/chai/ArrayManager.hpp | 2 +- tests/unit/array_manager_unit_tests.cpp | 12 ++++++------ 3 files changed, 8 insertions(+), 8 deletions(-) diff --git a/src/chai/ArrayManager.cpp b/src/chai/ArrayManager.cpp index a0ec5f23..c8778bbd 100644 --- a/src/chai/ArrayManager.cpp +++ b/src/chai/ArrayManager.cpp @@ -333,7 +333,7 @@ void ArrayManager::setUserCallback(void* pointer, UserCallback const& f) pointer_record->m_user_callback = f; } -void ArrayManager::setUserCallback(UserCallback const& f) +void ArrayManager::setGlobalUserCallback(UserCallback const& f) { m_user_callback = f; } diff --git a/src/chai/ArrayManager.hpp b/src/chai/ArrayManager.hpp index d0e96f8f..ccd88ff5 100644 --- a/src/chai/ArrayManager.hpp +++ b/src/chai/ArrayManager.hpp @@ -185,7 +185,7 @@ class ArrayManager * \brief Assign a user-defined callback triggered upon memory operations. * This callback applies to all ManagedArrays. */ - CHAISHAREDDLL_API void setUserCallback(UserCallback const& f); + CHAISHAREDDLL_API void setGlobalUserCallback(UserCallback const& f); /*! * \brief Set touched to false in all spaces for the given PointerRecord. diff --git a/tests/unit/array_manager_unit_tests.cpp b/tests/unit/array_manager_unit_tests.cpp index 5ee77714..1c212f21 100644 --- a/tests/unit/array_manager_unit_tests.cpp +++ b/tests/unit/array_manager_unit_tests.cpp @@ -144,9 +144,9 @@ TEST(ArrayManager, controlGlobalCallback) bool callbacksAreOn = false; // Set a global callback - arrayManager->setUserCallback([&] (const chai::PointerRecord*, chai::Action, chai::ExecutionSpace) { - callbacksAreOn = true; - }); + arrayManager->setGlobalUserCallback([&] (const chai::PointerRecord*, chai::Action, chai::ExecutionSpace) { + callbacksAreOn = true; + }); // Allocate an array and make sure the callback was called size_t sizeOfArray = 5; @@ -165,9 +165,9 @@ TEST(ArrayManager, controlGlobalCallback) // Now make sure the order doesn't matter for when the callback is set compared // to when callbacks are enabled - arrayManager->setUserCallback([&] (const chai::PointerRecord*, chai::Action, chai::ExecutionSpace) { - callbacksAreOn = true; - }); + arrayManager->setGlobalUserCallback([&] (const chai::PointerRecord*, chai::Action, chai::ExecutionSpace) { + callbacksAreOn = true; + }); // Reset the variable for testing if callbacks are on or off callbacksAreOn = false; From ee9dbaaad794afce6dbf090e5c858b0cf0dea559 Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Wed, 15 Apr 2020 08:40:23 -0700 Subject: [PATCH 42/70] address Alan`s comments. --- src/chai/ArrayManager.cpp | 15 +++++++------- src/chai/ArrayManager.hpp | 2 +- src/chai/ManagedArray.hpp | 4 ++-- src/chai/ManagedArray.inl | 36 +++++++--------------------------- src/chai/ManagedArray_thin.inl | 10 ++++++---- 5 files changed, 24 insertions(+), 43 deletions(-) diff --git a/src/chai/ArrayManager.cpp b/src/chai/ArrayManager.cpp index 89a04d3c..c3716553 100644 --- a/src/chai/ArrayManager.cpp +++ b/src/chai/ArrayManager.cpp @@ -172,18 +172,16 @@ ExecutionSpace ArrayManager::getExecutionSpace() void ArrayManager::registerTouch(PointerRecord* pointer_record) { if (m_current_execution_space == NONE) return; - if (pointer_record) { - registerTouch(pointer_record, m_current_execution_space); - } + registerTouch(pointer_record, m_current_execution_space); } void ArrayManager::registerTouch(PointerRecord* pointer_record, ExecutionSpace space) { - if (pointer_record) { - CHAI_LOG(Debug, pointer_record->m_pointers[space] << " touched in space " << space); + if (pointer_record && pointer_record != s_null_record) { if (space != NONE) { + CHAI_LOG(Debug, pointer_record->m_pointers[space] << " touched in space " << space); std::lock_guard lock(m_mutex); pointer_record->m_touched[space] = true; pointer_record->m_last_space = space; @@ -316,7 +314,6 @@ void ArrayManager::free(PointerRecord* pointer_record, ExecutionSpace spaceToFre } } - size_t ArrayManager::getSize(void* ptr) { // TODO @@ -365,7 +362,11 @@ PointerRecord* ArrayManager::makeManaged(void* pointer, auto pointer_record = getPointerRecord(pointer); if (pointer_record == &s_null_record) { - pointer_record = new PointerRecord(); + if (pointer) { + pointer_record = new PointerRecord(); + } else { + return pointer_record; + } } pointer_record->m_pointers[space] = pointer; pointer_record->m_owned[space] = owned; diff --git a/src/chai/ArrayManager.hpp b/src/chai/ArrayManager.hpp index 279d9fe7..6606075b 100644 --- a/src/chai/ArrayManager.hpp +++ b/src/chai/ArrayManager.hpp @@ -151,7 +151,7 @@ class ArrayManager /*! * \brief Free allocation(s) associated with the given PointerRecord. - * Default (space == NONE) will free all allocations and delete + * Default (space == NONE) will free all allocations and delete * the pointer record. */ CHAISHAREDDLL_API void free(PointerRecord* pointer, ExecutionSpace space = NONE); diff --git a/src/chai/ManagedArray.hpp b/src/chai/ManagedArray.hpp index f14c2576..bb9f01fd 100644 --- a/src/chai/ManagedArray.hpp +++ b/src/chai/ManagedArray.hpp @@ -112,7 +112,7 @@ class ManagedArray : public CHAICopyable */ CHAI_HOST_DEVICE ManagedArray(std::nullptr_t other); - CHAI_HOST_DEVICE ManagedArray(PointerRecord* record, ExecutionSpace space); + CHAI_HOST ManagedArray(PointerRecord* record, ExecutionSpace space); /*! * \brief Allocate data for the ManagedArray in the specified space. @@ -342,7 +342,7 @@ class ManagedArray : public CHAICopyable */ template ::value, typename std::enable_if::type = 0> - CHAI_HOST_DEVICE void moveInnerImpl(); + CHAI_HOST void moveInnerImpl(); #endif public: diff --git a/src/chai/ManagedArray.inl b/src/chai/ManagedArray.inl index 99be774e..27cba932 100644 --- a/src/chai/ManagedArray.inl +++ b/src/chai/ManagedArray.inl @@ -56,7 +56,7 @@ CHAI_HOST_DEVICE ManagedArray::ManagedArray( #if !defined(__CUDA_ARCH__) && !defined(__HIP_DEVICE_COMPILE__) this->allocate(elems, space); - + #if defined(CHAI_ENABLE_UM) if(space == UM) { m_pointer_record->m_pointers[CPU] = m_active_pointer; @@ -76,8 +76,6 @@ CHAI_HOST_DEVICE ManagedArray::ManagedArray( ManagedArray(spaces, allocators) { #if !defined(__CUDA_ARCH__) && !defined(__HIP_DEVICE_COMPILE__) - m_elems = elems; - this->allocate(elems, space); #if defined(CHAI_ENABLE_UM) @@ -92,23 +90,13 @@ CHAI_HOST_DEVICE ManagedArray::ManagedArray( template CHAI_INLINE CHAI_HOST_DEVICE ManagedArray::ManagedArray(std::nullptr_t) : - m_active_pointer(nullptr), - m_active_base_pointer(nullptr), - m_resource_manager(nullptr), - m_elems(0), - m_offset(0), - m_pointer_record(nullptr), - m_is_slice(false) + ManagedArray() { -#if !defined(__CUDA_ARCH__) && !defined(__HIP_DEVICE_COMPILE__) - m_resource_manager = ArrayManager::getInstance(); - m_pointer_record = &ArrayManager::s_null_record; -#endif } template CHAI_INLINE -CHAI_HOST_DEVICE ManagedArray::ManagedArray(PointerRecord* record, ExecutionSpace space): +CHAI_HOST ManagedArray::ManagedArray(PointerRecord* record, ExecutionSpace space): m_active_pointer(static_cast(record->m_pointers[space])), m_active_base_pointer(static_cast(record->m_pointers[space])), m_resource_manager(nullptr), @@ -117,12 +105,10 @@ CHAI_HOST_DEVICE ManagedArray::ManagedArray(PointerRecord* record, ExecutionS m_pointer_record(record), m_is_slice(false) { -#if !defined(__CUDA_ARCH__) && !defined(__HIP_DEVICE_COMPILE__) m_resource_manager = ArrayManager::getInstance(); if (m_pointer_record == nullptr) { m_pointer_record = &ArrayManager::s_null_record; } -#endif } @@ -164,7 +150,6 @@ CHAI_HOST_DEVICE ManagedArray::ManagedArray(T* data, ArrayManager* array_mana m_resource_manager = ArrayManager::getInstance(); } if (m_pointer_record == &ArrayManager::s_null_record || m_pointer_record==nullptr) { - bool owned = true; m_pointer_record = m_resource_manager->makeManaged((void *) data, sizeof(T)*m_elems,ExecutionSpace(CPU),true); } registerTouch(CPU); @@ -186,7 +171,7 @@ CHAI_HOST void ManagedArray::allocate( space = m_resource_manager->getDefaultAllocationSpace(); } if (m_pointer_record == &ArrayManager::s_null_record) { - m_pointer_record = m_resource_manager->makeManaged((void *) m_active_base_pointer,m_elems*sizeof(T),CPU,true); + m_pointer_record = m_resource_manager->makeManaged((void *) m_active_base_pointer,m_elems*sizeof(T),space,true); } m_pointer_record->m_user_callback = cback; @@ -203,7 +188,7 @@ CHAI_HOST void ManagedArray::allocate( // trigger a moveInnerImpl, which expects inner values to be initialized. initInner(); - CHAI_LOG(Debug, "m_active_ptr allocated at address: " << m_active_pointer); + CHAI_LOG(Debug, "m_active_base_ptr allocated at address: " << m_active_base_pointer); } } } @@ -266,7 +251,6 @@ CHAI_HOST void ManagedArray::free(ExecutionSpace space) if (space == NONE) { m_pointer_record = &ArrayManager::s_null_record; } - m_is_slice = false; } else { CHAI_LOG(Debug, "Cannot free a slice!"); } @@ -290,7 +274,7 @@ CHAI_INLINE CHAI_HOST void ManagedArray::registerTouch(ExecutionSpace space) { if (m_active_pointer && (m_pointer_record == nullptr || m_pointer_record == &ArrayManager::s_null_record)) { CHAI_LOG(Warning,"registerTouch called on ManagedArray with nullptr pointer record."); - m_pointer_record = m_resource_manager->makeManaged((void *)m_active_base_pointer,m_elems*sizeof(T),CPU,true); + m_pointer_record = m_resource_manager->makeManaged((void *)m_active_base_pointer,m_elems*sizeof(T),space,true); } m_resource_manager->registerTouch(m_pointer_record, space); } @@ -341,7 +325,7 @@ CHAI_HOST_DEVICE void ManagedArray::set(size_t i, T val) const { m_resource_manager->set(static_cast((void*)((char*)m_pointer_record->m_pointers[m_pointer_record->m_last_space]+sizeof(T)*m_offset)), i, val); #else m_active_pointer[i] = val; - #endif // !defined(__CUDA_ARCH__) + #endif // !defined(__CUDA_ARCH__) && !defined(__HIP_DEVICE_COMPILE__) } template @@ -385,9 +369,6 @@ CHAI_INLINE CHAI_HOST void ManagedArray::move(ExecutionSpace space) { - if (m_resource_manager == nullptr) { - m_resource_manager = ArrayManager::getInstance(); - } if (m_pointer_record != &ArrayManager::s_null_record) { CHAI_LOG(Debug, "Moving " << m_active_pointer); m_active_base_pointer = static_cast(m_resource_manager->move((void *)m_active_base_pointer, m_pointer_record, space)); @@ -414,9 +395,6 @@ CHAI_INLINE CHAI_HOST_DEVICE ManagedArray::operator T*() const { #if !defined(__CUDA_ARCH__) && !defined(__HIP_DEVICE_COMPILE__) if (m_active_pointer) { - if (m_resource_manager == nullptr) { - m_resource_manager = ArrayManager::getInstance(); - } if (m_pointer_record == nullptr || m_pointer_record == &ArrayManager::s_null_record) { CHAI_LOG(Warning, "nullptr pointer_record associated with non-nullptr active_pointer") } diff --git a/src/chai/ManagedArray_thin.inl b/src/chai/ManagedArray_thin.inl index 0655a990..06a704d5 100644 --- a/src/chai/ManagedArray_thin.inl +++ b/src/chai/ManagedArray_thin.inl @@ -158,8 +158,8 @@ CHAI_HOST void ManagedArray::reallocate(size_t new_elems) template CHAI_INLINE CHAI_HOST void ManagedArray::free(ExecutionSpace space) { - if (space == CPU || space == NONE) { - if (!m_is_slice) { + if (!m_is_slice) { + if (space == CPU || space == NONE) { #if defined(CHAI_ENABLE_UM) cudaFree(m_active_pointer); #else @@ -167,10 +167,12 @@ CHAI_INLINE CHAI_HOST void ManagedArray::free(ExecutionSpace space) #endif m_active_pointer = nullptr; m_active_base_pointer = nullptr; - } else { - CHAI_LOG(Debug, "tried to free slice!"); + m_elems = 0; } } + else { + CHAI_LOG(Debug, "tried to free slice!"); + } } From d0f3769b25f20b9030efcab2b9503d6f0be13579 Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Wed, 15 Apr 2020 09:05:18 -0700 Subject: [PATCH 43/70] fix gpu tests, add reportLeaks API. --- src/chai/ArrayManager.cpp | 49 ++++++++---- src/chai/ArrayManager.hpp | 11 ++- src/chai/ManagedArray.hpp | 1 + src/chai/ManagedArray.inl | 14 +++- src/chai/ManagedArray_thin.inl | 13 ++- src/chai/Types.hpp | 2 +- src/chai/managed_ptr.hpp | 2 +- src/util/forall.hpp | 2 + tests/integration/managed_array_tests.cpp | 96 +++++++++++++++++++---- 9 files changed, 149 insertions(+), 41 deletions(-) diff --git a/src/chai/ArrayManager.cpp b/src/chai/ArrayManager.cpp index c3716553..c7c65bd4 100644 --- a/src/chai/ArrayManager.cpp +++ b/src/chai/ArrayManager.cpp @@ -67,17 +67,21 @@ void ArrayManager::registerPointer( PointerRecord ** found_pointer_record_addr = found_pointer_record_pair->second; if (found_pointer_record_addr != nullptr) { - CHAI_LOG(Warning, "ArrayManager::registerPointer found a record for " << - pointer << " already there. Deleting abandoned pointer record."); - PointerRecord *foundRecord = *(found_pointer_record_pair->second); - callback(foundRecord, ACTION_FOUND_ABANDONED, space); + // if it's actually the same pointer record, then we're OK. If it's a different + // one, delete the old one. + if (foundRecord != record) { + CHAI_LOG(Warning, "ArrayManager::registerPointer found a record for " << + pointer << " already there. Deleting abandoned pointer record."); - for (int fspace = 0; fspace < NUM_EXECUTION_SPACES; ++fspace) { - foundRecord->m_pointers[fspace] = nullptr; - } + callback(foundRecord, ACTION_FOUND_ABANDONED, space); - delete foundRecord; + for (int fspace = 0; fspace < NUM_EXECUTION_SPACES; ++fspace) { + foundRecord->m_pointers[fspace] = nullptr; + } + + delete foundRecord; + } } } @@ -118,6 +122,7 @@ void ArrayManager::deregisterPointer(PointerRecord* record, bool deregisterFromU if (deregisterFromUmpire) { m_resource_manager.deregisterAllocation(pointer); } + CHAI_LOG(Debug, "DeRegistering " << pointer); m_pointer_map.erase(pointer); } } @@ -262,17 +267,13 @@ void ArrayManager::free(PointerRecord* pointer_record, ExecutionSpace spaceToFre for (int space = CPU; space < NUM_EXECUTION_SPACES; ++space) { if (space == spaceToFree || spaceToFree == NONE) { if (pointer_record->m_pointers[space]) { + void* space_ptr = pointer_record->m_pointers[space]; if (pointer_record->m_owned[space]) { - void* space_ptr = pointer_record->m_pointers[space]; #if defined(CHAI_ENABLE_UM) if (space_ptr == pointer_record->m_pointers[UM]) { callback(pointer_record, ACTION_FREE, ExecutionSpace(UM)); - { - std::lock_guard lock(m_mutex); - m_pointer_map.erase(space_ptr); - } auto alloc = m_resource_manager.getAllocator( pointer_record->m_allocators[space]); @@ -287,10 +288,6 @@ void ArrayManager::free(PointerRecord* pointer_record, ExecutionSpace spaceToFre callback(pointer_record, ACTION_FREE, ExecutionSpace(space)); - { - std::lock_guard lock(m_mutex); - m_pointer_map.erase(space_ptr); - } auto alloc = m_resource_manager.getAllocator( pointer_record->m_allocators[space]); @@ -305,6 +302,11 @@ void ArrayManager::free(PointerRecord* pointer_record, ExecutionSpace spaceToFre { m_resource_manager.deregisterAllocation(pointer_record->m_pointers[space]); } + { + std::lock_guard lock(m_mutex); + CHAI_LOG(Debug, "DeRegistering " << space_ptr); + m_pointer_map.erase(space_ptr); + } } } } @@ -361,6 +363,7 @@ PointerRecord* ArrayManager::makeManaged(void* pointer, {pointer, size, m_allocators[space]->getAllocationStrategy()}); auto pointer_record = getPointerRecord(pointer); + if (pointer_record == &s_null_record) { if (pointer) { pointer_record = new PointerRecord(); @@ -368,6 +371,11 @@ PointerRecord* ArrayManager::makeManaged(void* pointer, return pointer_record; } } + else { + CHAI_LOG(Warning, "ArrayManager::makeManaged found abandoned pointer record!!!"); + callback(pointer_record, ACTION_FOUND_ABANDONED, space); + } + pointer_record->m_pointers[space] = pointer; pointer_record->m_owned[space] = owned; pointer_record->m_size = size; @@ -447,6 +455,13 @@ size_t ArrayManager::getTotalSize() const return total; } +void ArrayManager::reportLeaks() const +{ + for (auto entry : m_pointer_map) { + callback(*entry.second, ACTION_LEAKED, NONE); + } +} + int ArrayManager::getAllocatorId(ExecutionSpace space) const { diff --git a/src/chai/ArrayManager.hpp b/src/chai/ArrayManager.hpp index 6606075b..db16063b 100644 --- a/src/chai/ArrayManager.hpp +++ b/src/chai/ArrayManager.hpp @@ -237,11 +237,20 @@ class ArrayManager */ CHAISHAREDDLL_API size_t getTotalSize() const; + /*! + * \brief Calls callbacks of pointers still in the map with ACTION_LEAKED. + */ + CHAISHAREDDLL_API void reportLeaks() const; + + /*! + * \brief Get the allocator ID + * + * \return The allocator ID. + */ CHAISHAREDDLL_API int getAllocatorId(ExecutionSpace space) const; /*! * \brief Wraps our resource manager's copy. - * */ CHAISHAREDDLL_API void copy(void * dst, void * src, size_t size); diff --git a/src/chai/ManagedArray.hpp b/src/chai/ManagedArray.hpp index bb9f01fd..fc1a1610 100644 --- a/src/chai/ManagedArray.hpp +++ b/src/chai/ManagedArray.hpp @@ -129,6 +129,7 @@ class ManagedArray : public CHAICopyable [] (const PointerRecord*, Action, ExecutionSpace) {}); + /*! * \brief Reallocate data for the ManagedArray. * diff --git a/src/chai/ManagedArray.inl b/src/chai/ManagedArray.inl index 27cba932..fd20bc29 100644 --- a/src/chai/ManagedArray.inl +++ b/src/chai/ManagedArray.inl @@ -126,10 +126,18 @@ CHAI_HOST_DEVICE ManagedArray::ManagedArray(ManagedArray const& other): #if !defined(__CUDA_ARCH__) && !defined(__HIP_DEVICE_COMPILE__) m_elems = other.m_pointer_record->m_size/sizeof(T); if (m_active_base_pointer) { - /// Move nested ManagedArrays first, so they are working with a valid m_active_pointer for the host, - // and so the meta data associated with them are updated before we move that down. - moveInnerImpl(); + ExecutionSpace prev_space = m_pointer_record->m_last_space; + if (prev_space == CPU) { + /// Move nested ManagedArrays first, so they are working with a valid m_active_pointer for the host, + // and so the meta data associated with them are updated before we move the other array down. + moveInnerImpl(); + } move(); + if (prev_space == GPU) { + /// Move nested ManagedArrays after the move, so they are working with a valid m_active_pointer for the host, + // and so the meta data associated with them are updated with live GPU data + moveInnerImpl(); + } } #endif } diff --git a/src/chai/ManagedArray_thin.inl b/src/chai/ManagedArray_thin.inl index 06a704d5..46307396 100644 --- a/src/chai/ManagedArray_thin.inl +++ b/src/chai/ManagedArray_thin.inl @@ -56,9 +56,15 @@ CHAI_HOST_DEVICE ManagedArray::ManagedArray(size_t elems, ExecutionSpace spac } -template -CHAI_INLINE -CHAI_HOST_DEVICE ManagedArray::ManagedArray(std::nullptr_t) +template +CHAI_INLINE CHAI_HOST_DEVICE ManagedArray::ManagedArray(std::nullptr_t) + : m_active_pointer(nullptr), + m_active_base_pointer(nullptr), + m_resource_manager(nullptr), + m_elems(0), + m_offset(0), + m_pointer_record(nullptr), + m_is_slice(false) { } @@ -260,6 +266,7 @@ ManagedArray::ManagedArray(T* data, CHAIDISAMBIGUATE, bool) : m_resource_manager(nullptr), m_elems(-1), m_pointer_record(nullptr), + m_offset(0), m_is_slice(false) { } diff --git a/src/chai/Types.hpp b/src/chai/Types.hpp index d6a017ba..aeb16bf7 100644 --- a/src/chai/Types.hpp +++ b/src/chai/Types.hpp @@ -29,7 +29,7 @@ namespace chai typedef unsigned int uint; - enum Action { ACTION_ALLOC, ACTION_FREE, ACTION_MOVE, ACTION_CAPTURED, ACTION_FOUND_ABANDONED }; + enum Action { ACTION_ALLOC, ACTION_FREE, ACTION_MOVE, ACTION_CAPTURED, ACTION_FOUND_ABANDONED, ACTION_LEAKED }; using UserCallback = std::function; } // end of namespace chai diff --git a/src/chai/managed_ptr.hpp b/src/chai/managed_ptr.hpp index 49ffa0e8..39dca4b8 100644 --- a/src/chai/managed_ptr.hpp +++ b/src/chai/managed_ptr.hpp @@ -173,7 +173,7 @@ namespace chai { /// /// Default constructor. /// - CHAI_HOST_DEVICE constexpr managed_ptr() noexcept = default; + constexpr managed_ptr() noexcept = default; /// /// @author Alan Dayton diff --git a/src/util/forall.hpp b/src/util/forall.hpp index f1ae1835..45e84e19 100644 --- a/src/util/forall.hpp +++ b/src/util/forall.hpp @@ -80,6 +80,8 @@ void forall(gpu, int begin, int end, LOOP_BODY&& body) hipLaunchKernelGGL(forall_kernel_gpu, dim3(gridSize), dim3(blockSize), 0,0, begin, end - begin, body); hipDeviceSynchronize(); +#elif defined(CHAI_ENABLE_GPU_SIMULATION_MODE) + forall_kernel_cpu(begin, end, body); #endif rm->setExecutionSpace(chai::NONE); diff --git a/tests/integration/managed_array_tests.cpp b/tests/integration/managed_array_tests.cpp index d87c94e6..403c34bc 100644 --- a/tests/integration/managed_array_tests.cpp +++ b/tests/integration/managed_array_tests.cpp @@ -5,7 +5,6 @@ // 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(); } \ @@ -17,6 +16,8 @@ #define device_assert(EXP) assert(EXP) #endif +#define assert_empty_map(IGNORED) ASSERT_EQ(chai::ArrayManager::getInstance()->getPointerMap().size(),0) + #include "chai/config.hpp" #include "../src/util/forall.hpp" @@ -38,6 +39,8 @@ TEST(ManagedArray, SetOnHost) forall(sequential(), 0, 10, [=](int i) { ASSERT_EQ(array[i], i); }); array.free(); + + assert_empty_map(true); } #if (!defined(CHAI_DISABLE_RM)) @@ -52,6 +55,7 @@ TEST(ManagedArray, Const) forall(sequential(), 0, 10, [=](int i) { ASSERT_EQ(array_const[i], i); }); array.free(); + assert_empty_map(true); } TEST(ManagedArray, UserCallbackHost) @@ -79,6 +83,7 @@ TEST(ManagedArray, Slice) { chai::ManagedArray sl = array.slice(0,5); sl.free(); array.free(); + assert_empty_map(true); } TEST(ManagedArray, SliceOfSlice) { @@ -102,6 +107,7 @@ TEST(ManagedArray, SliceOfSlice) { sl1.free(); sl2.free(); array.free(); + assert_empty_map(true); } #if defined(CHAI_ENABLE_PICK) @@ -117,6 +123,7 @@ TEST(ManagedArray, PickHostFromHostConst) { ASSERT_EQ(temp, 5); array.free(); + assert_empty_map(true); } #endif @@ -130,6 +137,7 @@ TEST(ManagedArray, PickHostFromHost) ASSERT_EQ(temp, 5); array.free(); + assert_empty_map(true); } TEST(ManagedArray, SetHostToHost) @@ -143,6 +151,7 @@ TEST(ManagedArray, SetHostToHost) ASSERT_EQ(array[5], 10); array.free(); + assert_empty_map(true); } @@ -168,6 +177,7 @@ TEST(ManagedArray, IncrementDecrementOnHost) arrayI.free(); arrayD.free(); + assert_empty_map(true); } @@ -184,7 +194,7 @@ TEST(ManagedArray, PickHostFromHostConstUM) { ASSERT_EQ(temp, 5); array.free(); - // array_const.free(); + assert_empty_map(true); } #endif @@ -198,6 +208,7 @@ TEST(ManagedArray, PickHostFromHostUM) ASSERT_EQ(temp, 5); array.free(); + assert_empty_map(true); } TEST(ManagedArray, SetHostToHostUM) @@ -211,6 +222,7 @@ TEST(ManagedArray, SetHostToHostUM) ASSERT_EQ(array[5], 10); array.free(); + assert_empty_map(true); } TEST(ManagedArray, IncrementDecrementOnHostUM) @@ -235,6 +247,7 @@ TEST(ManagedArray, IncrementDecrementOnHostUM) arrayI.free(); arrayD.free(); + assert_empty_map(true); } #endif @@ -260,6 +273,7 @@ GPU_TEST(ManagedArray, PickandSetDeviceToDeviceUM) array1.free(); array2.free(); + assert_empty_map(true); } GPU_TEST(ManagedArray, PickHostFromDeviceUM) @@ -272,6 +286,7 @@ GPU_TEST(ManagedArray, PickHostFromDeviceUM) ASSERT_EQ(temp, 5); array.free(); + assert_empty_map(true); } #if (!defined(CHAI_DISABLE_RM)) @@ -286,7 +301,7 @@ GPU_TEST(ManagedArray, PickHostFromDeviceConstUM) { ASSERT_EQ(temp, 5); array.free(); - // array_const.free(); + assert_empty_map(true); } #endif @@ -302,6 +317,7 @@ GPU_TEST(ManagedArray, SetHostToDeviceUM) ASSERT_EQ(temp, 10); array.free(); + assert_empty_map(true); } GPU_TEST(ManagedArray, IncrementDecrementOnDeviceUM) @@ -326,6 +342,7 @@ GPU_TEST(ManagedArray, IncrementDecrementOnDeviceUM) arrayI.free(); arrayD.free(); + assert_empty_map(true); } GPU_TEST(ManagedArray, IncrementDecrementFromHostOnDeviceUM) @@ -345,6 +362,7 @@ GPU_TEST(ManagedArray, IncrementDecrementFromHostOnDeviceUM) ASSERT_EQ(temp, 8); array.free(); + assert_empty_map(true); } GPU_TEST(ManagedArray, PickandSetSliceDeviceToDeviceUM) { @@ -367,6 +385,7 @@ GPU_TEST(ManagedArray, PickandSetSliceDeviceToDeviceUM) { }); array.free(); + assert_empty_map(true); } #endif @@ -390,6 +409,7 @@ GPU_TEST(ManagedArray, PickandSetDeviceToDevice) array1.free(); array2.free(); + assert_empty_map(true); } GPU_TEST(ManagedArray, PickandSetSliceDeviceToDevice) { @@ -412,6 +432,7 @@ GPU_TEST(ManagedArray, PickandSetSliceDeviceToDevice) { }); array.free(); + assert_empty_map(true); } @@ -425,6 +446,7 @@ GPU_TEST(ManagedArray, PickHostFromDevice) ASSERT_EQ(temp, 5); array.free(); + assert_empty_map(true); } GPU_TEST(ManagedArray, PickHostFromDeviceConst) @@ -441,7 +463,7 @@ GPU_TEST(ManagedArray, PickHostFromDeviceConst) ASSERT_EQ(temp, 5); array.free(); - // array_const.free(); + assert_empty_map(true); } GPU_TEST(ManagedArray, SetHostToDevice) @@ -456,6 +478,7 @@ GPU_TEST(ManagedArray, SetHostToDevice) ASSERT_EQ(temp, 10); array.free(); + assert_empty_map(true); } GPU_TEST(ManagedArray, IncrementDecrementOnDevice) { @@ -479,6 +502,7 @@ GPU_TEST(ManagedArray, IncrementDecrementOnDevice) arrayI.free(); arrayD.free(); + assert_empty_map(true); } GPU_TEST(ManagedArray, IncrementDecrementFromHostOnDevice) @@ -498,6 +522,7 @@ GPU_TEST(ManagedArray, IncrementDecrementFromHostOnDevice) ASSERT_EQ(temp, 8); array.free(); + assert_empty_map(true); } #endif #endif @@ -523,6 +548,7 @@ GPU_TEST(ManagedArray, SliceOfSliceDevice) { sl1.free(); sl2.free(); array.free(); + assert_empty_map(true); } GPU_TEST(ManagedArray, SliceDevice) { @@ -546,6 +572,7 @@ GPU_TEST(ManagedArray, SliceDevice) { sl1.free(); sl2.free(); array.free(); + assert_empty_map(true); } GPU_TEST(ManagedArray, SetOnDevice) { @@ -558,6 +585,7 @@ GPU_TEST(ManagedArray, SetOnDevice) { forall(sequential(), 0, 10, [=](int i) { ASSERT_EQ(array[i], 2 * i); }); array.free(); + assert_empty_map(true); } GPU_TEST(ManagedArray, GetGpuOnHost) @@ -569,6 +597,7 @@ GPU_TEST(ManagedArray, GetGpuOnHost) forall(sequential(), 0, 10, [=](int i) { ASSERT_EQ(array[i], i); }); array.free(); + assert_empty_map(true); } #if defined(CHAI_ENABLE_UM) @@ -581,6 +610,7 @@ GPU_TEST(ManagedArray, SetOnDeviceUM) forall(sequential(), 0, 10, [=](int i) { ASSERT_EQ(array[i], i); }); array.free(); + assert_empty_map(true); } #endif #endif @@ -594,6 +624,7 @@ TEST(ManagedArray, Allocate) ASSERT_EQ(array.size(), 10u); array.free(); + assert_empty_map(true); } TEST(ManagedArray, ReallocateCPU) @@ -616,6 +647,7 @@ TEST(ManagedArray, ReallocateCPU) }); array.free(); + assert_empty_map(true); } #if defined(CHAI_ENABLE_CUDA) || defined(CHAI_ENABLE_HIP) @@ -639,6 +671,7 @@ GPU_TEST(ManagedArray, ReallocateGPU) }); array.free(); + assert_empty_map(true); } #endif @@ -658,6 +691,7 @@ TEST(ManagedArray, NullpointerConversions) chai::ManagedArray c(nullptr); ASSERT_EQ(c.size(), 0u); + assert_empty_map(true); } #if defined(CHAI_ENABLE_IMPLICIT_CONVERSIONS) @@ -671,6 +705,7 @@ TEST(ManagedArray, ImplicitConversions) a.free(); SUCCEED(); + assert_empty_map(true); } #endif @@ -689,6 +724,7 @@ TEST(ManagedArray, PodTest) }); array.free(); + assert_empty_map(true); } #if defined(CHAI_ENABLE_CUDA) || defined(CHAI_ENABLE_HIP) @@ -707,6 +743,7 @@ GPU_TEST(ManagedArray, PodTestGPU) }); array.free(); + assert_empty_map(true); } #endif @@ -731,6 +768,7 @@ TEST(ManagedArray, ExternalConstructorUnowned) } std::free(data); + assert_empty_map(true); } TEST(ManagedArray, ExternalConstructorOwned) @@ -747,6 +785,7 @@ TEST(ManagedArray, ExternalConstructorOwned) forall(sequential(), 0, 20, [=](int i) { ASSERT_EQ(data[i], array[i]); }); array.free(); + assert_empty_map(true); } #if defined(CHAI_ENABLE_CUDA) || defined(CHAI_ENABLE_HIP) GPU_TEST(ManagedArray, ExternalUnownedMoveToGPU) @@ -764,6 +803,7 @@ GPU_TEST(ManagedArray, ExternalUnownedMoveToGPU) forall(sequential(), 0, 20, [=] (int i) { ASSERT_EQ(array[i], 1.0f * i); }); array.free(); + assert_empty_map(true); } #endif #endif @@ -776,6 +816,7 @@ TEST(ManagedArray, Reset) array.reset(); array.free(); + assert_empty_map(true); } #if defined(CHAI_ENABLE_CUDA) || defined(CHAI_ENABLE_HIP) @@ -793,6 +834,7 @@ GPU_TEST(ManagedArray, ResetDevice) forall(sequential(), 0, 20, [=](int i) { ASSERT_EQ(array[i], 0.0f); }); array.free(); + assert_empty_map(true); } #endif #endif @@ -850,6 +892,7 @@ GPU_TEST(ManagedArray, UserCallback) ASSERT_EQ(bytes_alloc, 2 * 20 * sizeof(float)); ASSERT_EQ(bytes_free, 2 * 20 * sizeof(float)); + assert_empty_map(true); } GPU_TEST(ManagedArray, CallBackConst) @@ -874,6 +917,10 @@ GPU_TEST(ManagedArray, CallBackConst) ++num_h2d; } } + if (act == chai::ACTION_FOUND_ABANDONED) { + printf("in abandoned!\n"); + ASSERT_EQ(false,true); + } }; chai::ManagedArray array(100); @@ -907,6 +954,7 @@ GPU_TEST(ManagedArray, CallBackConst) ASSERT_EQ(num_d2h, 0); array.free(); + assert_empty_map(true); } GPU_TEST(ManagedArray, CallBackConstArray) @@ -931,6 +979,12 @@ GPU_TEST(ManagedArray, CallBackConstArray) ++num_h2d; } } + if (act == chai::ACTION_FOUND_ABANDONED) { + printf("in abandoned!\n"); + while(true) { + + } + } }; const int N = 5; @@ -951,12 +1005,9 @@ GPU_TEST(ManagedArray, CallBackConstArray) chai::ManagedArray errorTemp(N); - forall(sequential(), 0, N, - [=](int j) - { - temp[j] = N * i + j; - } - ); + forall(sequential(), 0, N, [=](int j) { + temp[j] = N * i + j; + }); outerArray[i] = temp; outerErrorArray[i] = errorTemp; @@ -1007,6 +1058,7 @@ GPU_TEST(ManagedArray, CallBackConstArray) outerArray.free(); outerErrorArray.free(); + assert_empty_map(true); } GPU_TEST(ManagedArray, CallBackConstArrayConst) @@ -1108,6 +1160,7 @@ GPU_TEST(ManagedArray, CallBackConstArrayConst) outerArray.free(); outerErrorArray.free(); + assert_empty_map(true); } #endif @@ -1127,6 +1180,7 @@ GPU_TEST(ManagedArray, Move) ASSERT_EQ(array[5], 5); array.free(); + assert_empty_map(true); } /** @@ -1175,6 +1229,7 @@ GPU_TEST(ManagedArray, MoveInnerToHost) } outerArray.free(); + assert_empty_map(true); } /** @@ -1237,6 +1292,7 @@ GPU_TEST(ManagedArray, MoveInnerToDevice) } outerArray.free(); + assert_empty_map(true); } /** @@ -1315,6 +1371,7 @@ GPU_TEST(ManagedArray, MoveInnerToDevice2) outerArray[i].free(); } outerArray.free(); + assert_empty_map(true); } GPU_TEST(ManagedArray, MoveInnerToDeviceAgain) @@ -1382,6 +1439,7 @@ GPU_TEST(ManagedArray, MoveInnerToDeviceAgain) } outerArray.free(); + assert_empty_map(true); } #endif // CHAI_DISABLE_RM #endif // defined(CHAI_ENABLE_CUDA) || defined(CHAI_ENABLE_HIP) @@ -1407,6 +1465,7 @@ TEST(ManagedArray, DeepCopy) array.free(); copy.free(); + assert_empty_map(true); } #endif @@ -1431,18 +1490,18 @@ GPU_TEST(ManagedArray, DeviceDeepCopy) array.free(); copy.free(); + assert_empty_map(true); } -#endif -#endif // defined(CHAI_ENABLE_CUDA) || defined(CHAI_ENABLE_HIP) - -#ifdef CHAI_ENABLE_CUDA GPU_TEST(ManagedArray, CopyConstruct) { const int expectedValue = rand(); chai::ManagedArray array(1, chai::CPU); - array[0] = expectedValue; + + forall(sequential(), 0, 1, [=] (int i) { + array[i] = expectedValue; + }); chai::ManagedArray array2 = array; @@ -1454,8 +1513,13 @@ GPU_TEST(ManagedArray, CopyConstruct) results.move(chai::CPU); ASSERT_EQ(results[0], expectedValue); + + array.free(); + results.free(); + assert_empty_map(true); } +#endif #endif TEST(ManagedArray, SizeZero) @@ -1464,6 +1528,7 @@ TEST(ManagedArray, SizeZero) ASSERT_EQ(array.size(), 0u); array.allocate(0); ASSERT_EQ(array.size(), 0u); + assert_empty_map(true); } #if defined(CHAI_ENABLE_CUDA) || defined(CHAI_ENABLE_HIP) @@ -1482,5 +1547,6 @@ GPU_TEST(ManagedArray, CopyZero) }); array.free(); + assert_empty_map(true); } #endif From 7be427f006afb35ac51fe1cc63b24196dfddd407 Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Wed, 15 Apr 2020 09:12:13 -0700 Subject: [PATCH 44/70] bring moveInnerImpl into move command, centralize use of move a bit. --- src/chai/ManagedArray.inl | 28 +++++++++++++--------------- 1 file changed, 13 insertions(+), 15 deletions(-) diff --git a/src/chai/ManagedArray.inl b/src/chai/ManagedArray.inl index fd20bc29..651f6170 100644 --- a/src/chai/ManagedArray.inl +++ b/src/chai/ManagedArray.inl @@ -126,18 +126,7 @@ CHAI_HOST_DEVICE ManagedArray::ManagedArray(ManagedArray const& other): #if !defined(__CUDA_ARCH__) && !defined(__HIP_DEVICE_COMPILE__) m_elems = other.m_pointer_record->m_size/sizeof(T); if (m_active_base_pointer) { - ExecutionSpace prev_space = m_pointer_record->m_last_space; - if (prev_space == CPU) { - /// Move nested ManagedArrays first, so they are working with a valid m_active_pointer for the host, - // and so the meta data associated with them are updated before we move the other array down. - moveInnerImpl(); - } move(); - if (prev_space == GPU) { - /// Move nested ManagedArrays after the move, so they are working with a valid m_active_pointer for the host, - // and so the meta data associated with them are updated with live GPU data - moveInnerImpl(); - } } #endif } @@ -378,6 +367,12 @@ CHAI_HOST void ManagedArray::move(ExecutionSpace space) { if (m_pointer_record != &ArrayManager::s_null_record) { + ExecutionSpace prev_space = m_pointer_record->m_last_space; + if (prev_space == CPU) { + /// Move nested ManagedArrays first, so they are working with a valid m_active_pointer for the host, + // and so the meta data associated with them are updated before we move the other array down. + moveInnerImpl(); + } CHAI_LOG(Debug, "Moving " << m_active_pointer); m_active_base_pointer = static_cast(m_resource_manager->move((void *)m_active_base_pointer, m_pointer_record, space)); m_active_pointer = m_active_base_pointer + m_offset; @@ -387,6 +382,11 @@ void ManagedArray::move(ExecutionSpace space) CHAI_LOG(Debug, "T is non-const, registering touch of pointer" << m_active_pointer); m_resource_manager->registerTouch(m_pointer_record); } + if (prev_space == GPU) { + /// Move nested ManagedArrays after the move, so they are working with a valid m_active_pointer for the host, + // and so the meta data associated with them are updated with live GPU data + moveInnerImpl(); + } } } @@ -408,13 +408,11 @@ CHAI_HOST_DEVICE ManagedArray::operator T*() const { } ExecutionSpace prev_space = m_resource_manager->getExecutionSpace(); m_resource_manager->setExecutionSpace(CPU); - auto non_const_active_base_pointer = const_cast(static_cast(m_active_base_pointer)); - m_active_base_pointer = static_cast(m_resource_manager->move(non_const_active_base_pointer, m_pointer_record)); - m_active_pointer = m_active_base_pointer+m_offset; + move(); + // always touch regarless of constness of type (don't trust the application not to const-cast) m_resource_manager->registerTouch(m_pointer_record); - // Reset to whatever space we rode in on m_resource_manager->setExecutionSpace(prev_space); } From e034ace960542342caf4b4e6eb52c59f552cfec0 Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Wed, 15 Apr 2020 09:22:51 -0700 Subject: [PATCH 45/70] build fixes --- src/chai/ArrayManager.cpp | 2 +- src/chai/ManagedArray.hpp | 6 +++--- src/chai/ManagedArray.inl | 6 +++--- 3 files changed, 7 insertions(+), 7 deletions(-) diff --git a/src/chai/ArrayManager.cpp b/src/chai/ArrayManager.cpp index c7c65bd4..fb2cc932 100644 --- a/src/chai/ArrayManager.cpp +++ b/src/chai/ArrayManager.cpp @@ -183,7 +183,7 @@ void ArrayManager::registerTouch(PointerRecord* pointer_record) void ArrayManager::registerTouch(PointerRecord* pointer_record, ExecutionSpace space) { - if (pointer_record && pointer_record != s_null_record) { + if (pointer_record && pointer_record != &s_null_record) { if (space != NONE) { CHAI_LOG(Debug, pointer_record->m_pointers[space] << " touched in space " << space); diff --git a/src/chai/ManagedArray.hpp b/src/chai/ManagedArray.hpp index fc1a1610..964e92b7 100644 --- a/src/chai/ManagedArray.hpp +++ b/src/chai/ManagedArray.hpp @@ -166,7 +166,7 @@ class ManagedArray : public CHAICopyable */ CHAI_HOST void registerTouch(ExecutionSpace space); - CHAI_HOST void move(ExecutionSpace space=NONE); + CHAI_HOST void move(ExecutionSpace space=NONE) const; CHAI_HOST_DEVICE ManagedArray slice(size_t begin, size_t elems=(size_t)-1) const; @@ -330,7 +330,7 @@ class ManagedArray : public CHAICopyable */ template ::value, typename std::enable_if::type = 0> - CHAI_HOST void moveInnerImpl(); + CHAI_HOST void moveInnerImpl() const; /*! * \brief Does nothing since the inner data type does not inherit from @@ -343,7 +343,7 @@ class ManagedArray : public CHAICopyable */ template ::value, typename std::enable_if::type = 0> - CHAI_HOST void moveInnerImpl(); + CHAI_HOST void moveInnerImpl() const; #endif public: diff --git a/src/chai/ManagedArray.inl b/src/chai/ManagedArray.inl index 651f6170..bc300431 100644 --- a/src/chai/ManagedArray.inl +++ b/src/chai/ManagedArray.inl @@ -364,7 +364,7 @@ CHAI_HOST_DEVICE void ManagedArray::decr(size_t i) const { template CHAI_INLINE CHAI_HOST -void ManagedArray::move(ExecutionSpace space) +void ManagedArray::move(ExecutionSpace space) const { if (m_pointer_record != &ArrayManager::s_null_record) { ExecutionSpace prev_space = m_pointer_record->m_last_space; @@ -594,7 +594,7 @@ template::type> CHAI_INLINE CHAI_HOST void -ManagedArray::moveInnerImpl() +ManagedArray::moveInnerImpl() const { int len = m_pointer_record->m_size / sizeof(T); T * host_ptr = (T *) m_pointer_record->m_pointers[CPU]; @@ -612,7 +612,7 @@ template::type> CHAI_INLINE CHAI_HOST void -ManagedArray::moveInnerImpl() +ManagedArray::moveInnerImpl() const { } From 400350423327132ecfb60ad118ba18fee249942c Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Thu, 16 Apr 2020 10:03:21 -0700 Subject: [PATCH 46/70] fix segfaults in tests, use of initialized memory in managed_array_tests --- src/chai/ArrayManager.cpp | 4 ++++ src/chai/ManagedArray.inl | 3 ++- tests/integration/managed_array_tests.cpp | 3 ++- 3 files changed, 8 insertions(+), 2 deletions(-) diff --git a/src/chai/ArrayManager.cpp b/src/chai/ArrayManager.cpp index fb2cc932..de55d252 100644 --- a/src/chai/ArrayManager.cpp +++ b/src/chai/ArrayManager.cpp @@ -358,6 +358,10 @@ PointerRecord* ArrayManager::makeManaged(void* pointer, ExecutionSpace space, bool owned) { + if (space == NONE) { + space = getDefaultAllocationSpace(); + } + m_resource_manager.registerAllocation( pointer, {pointer, size, m_allocators[space]->getAllocationStrategy()}); diff --git a/src/chai/ManagedArray.inl b/src/chai/ManagedArray.inl index bc300431..7ad1c92f 100644 --- a/src/chai/ManagedArray.inl +++ b/src/chai/ManagedArray.inl @@ -168,7 +168,8 @@ CHAI_HOST void ManagedArray::allocate( space = m_resource_manager->getDefaultAllocationSpace(); } if (m_pointer_record == &ArrayManager::s_null_record) { - m_pointer_record = m_resource_manager->makeManaged((void *) m_active_base_pointer,m_elems*sizeof(T),space,true); + // since we are about to allocate, this will get registered + m_pointer_record = new PointerRecord(); } m_pointer_record->m_user_callback = cback; diff --git a/tests/integration/managed_array_tests.cpp b/tests/integration/managed_array_tests.cpp index 403c34bc..75749013 100644 --- a/tests/integration/managed_array_tests.cpp +++ b/tests/integration/managed_array_tests.cpp @@ -697,7 +697,8 @@ TEST(ManagedArray, NullpointerConversions) #if defined(CHAI_ENABLE_IMPLICIT_CONVERSIONS) TEST(ManagedArray, ImplicitConversions) { - chai::ManagedArray a(10); + chai::ManagedArray a(1); + a[0] = 3.14159; chai::ManagedArray a2 = a; From f7b84099664e1a77cf5285630224b990f8a7f03d Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Thu, 16 Apr 2020 11:11:11 -0700 Subject: [PATCH 47/70] address reviewer comments. --- src/chai/ArrayManager.cpp | 9 ++++----- src/util/forall.hpp | 6 +++--- tests/integration/managed_array_tests.cpp | 9 ++++++--- 3 files changed, 13 insertions(+), 11 deletions(-) diff --git a/src/chai/ArrayManager.cpp b/src/chai/ArrayManager.cpp index de55d252..d69fa87f 100644 --- a/src/chai/ArrayManager.cpp +++ b/src/chai/ArrayManager.cpp @@ -67,7 +67,7 @@ void ArrayManager::registerPointer( PointerRecord ** found_pointer_record_addr = found_pointer_record_pair->second; if (found_pointer_record_addr != nullptr) { - PointerRecord *foundRecord = *(found_pointer_record_pair->second); + PointerRecord *foundRecord = *found_pointer_record_addr; // if it's actually the same pointer record, then we're OK. If it's a different // one, delete the old one. if (foundRecord != record) { @@ -76,7 +76,7 @@ void ArrayManager::registerPointer( callback(foundRecord, ACTION_FOUND_ABANDONED, space); - for (int fspace = 0; fspace < NUM_EXECUTION_SPACES; ++fspace) { + for (int fspace = CPU; fspace < NUM_EXECUTION_SPACES; ++fspace) { foundRecord->m_pointers[fspace] = nullptr; } @@ -122,7 +122,7 @@ void ArrayManager::deregisterPointer(PointerRecord* record, bool deregisterFromU if (deregisterFromUmpire) { m_resource_manager.deregisterAllocation(pointer); } - CHAI_LOG(Debug, "DeRegistering " << pointer); + CHAI_LOG(Debug, "De-registering " << pointer); m_pointer_map.erase(pointer); } } @@ -176,7 +176,6 @@ ExecutionSpace ArrayManager::getExecutionSpace() void ArrayManager::registerTouch(PointerRecord* pointer_record) { - if (m_current_execution_space == NONE) return; registerTouch(pointer_record, m_current_execution_space); } @@ -300,7 +299,7 @@ void ArrayManager::free(PointerRecord* pointer_record, ExecutionSpace spaceToFre } else { - m_resource_manager.deregisterAllocation(pointer_record->m_pointers[space]); + m_resource_manager.deregisterAllocation(space_ptr); } { std::lock_guard lock(m_mutex); diff --git a/src/util/forall.hpp b/src/util/forall.hpp index 45e84e19..e992b0c5 100644 --- a/src/util/forall.hpp +++ b/src/util/forall.hpp @@ -73,15 +73,15 @@ void forall(gpu, int begin, int end, LOOP_BODY&& body) size_t blockSize = 32; size_t gridSize = (end - begin + blockSize - 1) / blockSize; -#if defined(CHAI_ENABLE_CUDA) +#if defined(CHAI_ENABLE_GPU_SIMULATION_MODE) + forall_kernel_cpu(begin, end, body); +#elif defined(CHAI_ENABLE_CUDA) forall_kernel_gpu<<>>(begin, end - begin, body); cudaDeviceSynchronize(); #elif defined(CHAI_ENABLE_HIP) hipLaunchKernelGGL(forall_kernel_gpu, dim3(gridSize), dim3(blockSize), 0,0, begin, end - begin, body); hipDeviceSynchronize(); -#elif defined(CHAI_ENABLE_GPU_SIMULATION_MODE) - forall_kernel_cpu(begin, end, body); #endif rm->setExecutionSpace(chai::NONE); diff --git a/tests/integration/managed_array_tests.cpp b/tests/integration/managed_array_tests.cpp index 75749013..8ff1be63 100644 --- a/tests/integration/managed_array_tests.cpp +++ b/tests/integration/managed_array_tests.cpp @@ -16,7 +16,12 @@ #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" @@ -982,9 +987,7 @@ GPU_TEST(ManagedArray, CallBackConstArray) } if (act == chai::ACTION_FOUND_ABANDONED) { printf("in abandoned!\n"); - while(true) { - - } + ASSERT_TRUE(false); } }; From 5d96a82d86967116ee419483da235e68be83ce4c Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Thu, 16 Apr 2020 11:26:01 -0700 Subject: [PATCH 48/70] fix thin mode build. --- src/chai/ManagedArray_thin.inl | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/src/chai/ManagedArray_thin.inl b/src/chai/ManagedArray_thin.inl index 46307396..77bb0d9e 100644 --- a/src/chai/ManagedArray_thin.inl +++ b/src/chai/ManagedArray_thin.inl @@ -169,7 +169,7 @@ CHAI_INLINE CHAI_HOST void ManagedArray::free(ExecutionSpace space) #if defined(CHAI_ENABLE_UM) cudaFree(m_active_pointer); #else - ::free(m_active_pointer); + ::free((void *)m_active_pointer); #endif m_active_pointer = nullptr; m_active_base_pointer = nullptr; @@ -239,7 +239,7 @@ CHAI_INLINE CHAI_HOST void ManagedArray::registerTouch(ExecutionSpace) } template -CHAI_INLINE CHAI_HOST void ManagedArray::move(ExecutionSpace) +CHAI_INLINE CHAI_HOST void ManagedArray::move(ExecutionSpace) const { } From 1080db7539262d73b23431172344a990d24ffea8 Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Thu, 16 Apr 2020 13:31:29 -0700 Subject: [PATCH 49/70] getInstance should be host, typo in docs. --- docs/sphinx/advanced_configuration.rst | 2 +- src/chai/ArrayManager.cpp | 1 - src/chai/ArrayManager.hpp | 1 - 3 files changed, 1 insertion(+), 3 deletions(-) diff --git a/docs/sphinx/advanced_configuration.rst b/docs/sphinx/advanced_configuration.rst index cec6e697..bbf4bfea 100644 --- a/docs/sphinx/advanced_configuration.rst +++ b/docs/sphinx/advanced_configuration.rst @@ -37,7 +37,7 @@ These arguments are explained in more detail below: GPU_SIMULATION_MODE support, then only the ``CPU`` execution space is available for use. * ENABLE_GPU_SIMULATION_MODE - This option simulates GPU support by enableing the GPU execution space, backed by a HOST + This option simulates GPU support by enabling the GPU execution space, backed by a HOST umpire allocator. If CHAI is built without CUDA, HIP, or GPU_SIMULATION_MODE support, then only the ``CPU`` execution space is available for use. diff --git a/src/chai/ArrayManager.cpp b/src/chai/ArrayManager.cpp index d69fa87f..cd6b7950 100644 --- a/src/chai/ArrayManager.cpp +++ b/src/chai/ArrayManager.cpp @@ -15,7 +15,6 @@ namespace chai PointerRecord ArrayManager::s_null_record = PointerRecord(); -CHAI_HOST_DEVICE ArrayManager* ArrayManager::getInstance() { static ArrayManager s_resource_manager_instance; diff --git a/src/chai/ArrayManager.hpp b/src/chai/ArrayManager.hpp index db16063b..c4b16b65 100644 --- a/src/chai/ArrayManager.hpp +++ b/src/chai/ArrayManager.hpp @@ -61,7 +61,6 @@ class ArrayManager * */ CHAISHAREDDLL_API - CHAI_HOST_DEVICE static ArrayManager* getInstance(); /*! From 176166f627d7fb2d64464ca483dc59ec1bcb1a4e Mon Sep 17 00:00:00 2001 From: Alan Dayton Date: Sat, 18 Apr 2020 12:00:06 -0700 Subject: [PATCH 50/70] Add method to get allocator --- src/chai/ArrayManager.hpp | 7 +++++++ src/chai/ArrayManager.inl | 5 +++++ 2 files changed, 12 insertions(+) diff --git a/src/chai/ArrayManager.hpp b/src/chai/ArrayManager.hpp index ccd88ff5..27b92fc9 100644 --- a/src/chai/ArrayManager.hpp +++ b/src/chai/ArrayManager.hpp @@ -311,6 +311,13 @@ class ArrayManager */ CHAISHAREDDLL_API void * frontOfAllocation(void * pointer); + /*! + * \brief Get the allocator for an execution space. + * + * \param space Execution space of the allocator to get. + * \param allocator The allocator for this space. + */ + void getAllocator(ExecutionSpace space, umpire::Allocator &allocator); /*! * \brief set the allocator for an execution space. diff --git a/src/chai/ArrayManager.inl b/src/chai/ArrayManager.inl index d7b0b214..d6121c24 100644 --- a/src/chai/ArrayManager.inl +++ b/src/chai/ArrayManager.inl @@ -103,6 +103,11 @@ void ArrayManager::copy(void * dst, void * src, size_t size) { m_resource_manager.copy(dst,src,size); } +CHAI_INLINE +void ArrayManager::getAllocator(ExecutionSpace space, umpire::Allocator &allocator) { + allocator = *m_allocators[space]; +} + CHAI_INLINE void ArrayManager::setAllocator(ExecutionSpace space, umpire::Allocator &allocator) { *m_allocators[space] = allocator; From 627beb51619598289eac6e8ac63e28f10ebe8a4f Mon Sep 17 00:00:00 2001 From: Alan Dayton Date: Mon, 20 Apr 2020 09:25:24 -0700 Subject: [PATCH 51/70] Return allocator by value --- src/chai/ArrayManager.hpp | 5 +++-- src/chai/ArrayManager.inl | 4 ++-- 2 files changed, 5 insertions(+), 4 deletions(-) diff --git a/src/chai/ArrayManager.hpp b/src/chai/ArrayManager.hpp index 27b92fc9..ea7ffcb4 100644 --- a/src/chai/ArrayManager.hpp +++ b/src/chai/ArrayManager.hpp @@ -315,9 +315,10 @@ class ArrayManager * \brief Get the allocator for an execution space. * * \param space Execution space of the allocator to get. - * \param allocator The allocator for this space. + * + * \return The allocator for the given space. */ - void getAllocator(ExecutionSpace space, umpire::Allocator &allocator); + umpire::Allocator getAllocator(ExecutionSpace space); /*! * \brief set the allocator for an execution space. diff --git a/src/chai/ArrayManager.inl b/src/chai/ArrayManager.inl index d6121c24..7ac78a63 100644 --- a/src/chai/ArrayManager.inl +++ b/src/chai/ArrayManager.inl @@ -104,8 +104,8 @@ void ArrayManager::copy(void * dst, void * src, size_t size) { } CHAI_INLINE -void ArrayManager::getAllocator(ExecutionSpace space, umpire::Allocator &allocator) { - allocator = *m_allocators[space]; +umpire::Allocator ArrayManager::getAllocator(ExecutionSpace space) { + return *m_allocators[space]; } CHAI_INLINE From 43695443dfefb1104c238fb41e7f3d899f5be064 Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Tue, 21 Apr 2020 08:58:48 -0700 Subject: [PATCH 52/70] cleanup old comment, initialize allocators when instantiating PointerRecord in allocate. --- src/chai/ArrayManager.cpp | 7 ------- src/chai/ManagedArray.inl | 7 +++++-- 2 files changed, 5 insertions(+), 9 deletions(-) diff --git a/src/chai/ArrayManager.cpp b/src/chai/ArrayManager.cpp index cd6b7950..eccb0d66 100644 --- a/src/chai/ArrayManager.cpp +++ b/src/chai/ArrayManager.cpp @@ -391,13 +391,6 @@ PointerRecord* ArrayManager::makeManaged(void* pointer, registerPointer(pointer_record, space, owned); } - // TODO Is this a problem? - // for (int i = 0; i < NUM_EXECUTION_SPACES; i++) { - // // If pointer is already active on some execution space, return that - // pointer if(pointer_record->m_touched[i] == true) - // return pointer_record->m_pointers[i]; - // } - return pointer_record; } diff --git a/src/chai/ManagedArray.inl b/src/chai/ManagedArray.inl index 7ad1c92f..47795607 100644 --- a/src/chai/ManagedArray.inl +++ b/src/chai/ManagedArray.inl @@ -168,8 +168,11 @@ CHAI_HOST void ManagedArray::allocate( space = m_resource_manager->getDefaultAllocationSpace(); } if (m_pointer_record == &ArrayManager::s_null_record) { - // since we are about to allocate, this will get registered - m_pointer_record = new PointerRecord(); + // since we are about to allocate, this will get registered + m_pointer_record = new PointerRecord(); + for (int space = CPU; space < NUM_EXECUTION_SPACES; ++space) { + m_pointer_record->m_allocators[space] = m_resource_manager->getAllocatorId(ExecutionSpace(space)); + } } m_pointer_record->m_user_callback = cback; From ea1898f08fec0d10b298b32697d522ba2e52d9a2 Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Tue, 21 Apr 2020 15:49:43 -0700 Subject: [PATCH 53/70] initializing resource_manager for nullptr assignment, make m_elems update in copy constructor dependent on non-null pointer_record. (pointer record can be nullptr for inner arrays device side initialized to nullptr --- src/chai/ManagedArray.inl | 9 +++++++-- 1 file changed, 7 insertions(+), 2 deletions(-) diff --git a/src/chai/ManagedArray.inl b/src/chai/ManagedArray.inl index 47795607..08db5cf1 100644 --- a/src/chai/ManagedArray.inl +++ b/src/chai/ManagedArray.inl @@ -124,8 +124,11 @@ CHAI_HOST_DEVICE ManagedArray::ManagedArray(ManagedArray const& other): m_is_slice(other.m_is_slice) { #if !defined(__CUDA_ARCH__) && !defined(__HIP_DEVICE_COMPILE__) - m_elems = other.m_pointer_record->m_size/sizeof(T); if (m_active_base_pointer) { + // we only update m_elems if we are not null and we have a pointer record + if (m_pointer_record) { + m_elems = m_pointer_record->m_size/sizeof(T); + } move(); } #endif @@ -372,7 +375,7 @@ void ManagedArray::move(ExecutionSpace space) const { if (m_pointer_record != &ArrayManager::s_null_record) { ExecutionSpace prev_space = m_pointer_record->m_last_space; - if (prev_space == CPU) { + if (prev_space == CPU || prev_space == NONE) { /// Move nested ManagedArrays first, so they are working with a valid m_active_pointer for the host, // and so the meta data associated with them are updated before we move the other array down. moveInnerImpl(); @@ -529,8 +532,10 @@ ManagedArray::operator= (std::nullptr_t) { m_offset = 0; #if !defined(__CUDA_ARCH__) && !defined(__HIP_DEVICE_COMPILE__) m_pointer_record = &ArrayManager::s_null_record; + m_resource_manager = ArrayManager::getInstance(); #else m_pointer_record = nullptr; + m_resource_manager = nullptr; #endif m_is_slice = false; return *this; From ad88283af20e1a7788275ac724e2e55d8b1a2af6 Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Tue, 21 Apr 2020 18:14:13 -0700 Subject: [PATCH 54/70] add unit tests for device initialized inner arrays, fix forall to get them to pass. --- src/chai/ManagedArray.hpp | 6 ++- src/chai/ManagedArray.inl | 10 +++- src/util/forall.hpp | 4 +- tests/integration/managed_array_tests.cpp | 59 +++++++++++++++++++++++ 4 files changed, 74 insertions(+), 5 deletions(-) diff --git a/src/chai/ManagedArray.hpp b/src/chai/ManagedArray.hpp index 964e92b7..ef21b3c1 100644 --- a/src/chai/ManagedArray.hpp +++ b/src/chai/ManagedArray.hpp @@ -378,18 +378,20 @@ class ManagedArray : public CHAICopyable // trigger a moveInnerImpl, which expects inner values to be initialized. template ::value, typename std::enable_if::type = 0> - CHAI_HOST void initInner(size_t start = 0) + CHAI_HOST bool initInner(size_t start = 0) { for (size_t i = start; i < m_elems; ++i) { m_active_base_pointer[i] = nullptr; } + return true; } // Do not deep initialize if T is not a CHAICopyable. template ::value, typename std::enable_if::type = 0> - CHAI_HOST void initInner(size_t = 0) + CHAI_HOST bool initInner(size_t = 0) { + return false; } #endif protected: diff --git a/src/chai/ManagedArray.inl b/src/chai/ManagedArray.inl index 08db5cf1..d9fc3ba2 100644 --- a/src/chai/ManagedArray.inl +++ b/src/chai/ManagedArray.inl @@ -224,7 +224,15 @@ CHAI_HOST void ManagedArray::reallocate(size_t elems) // if T is a CHAICopyable, then it is important to initialize all the new // ManagedArrays to nullptr at allocation, since it is extremely easy to // trigger a moveInnerImpl, which expects inner values to be initialized. - initInner(old_size); + if (initInner(old_size)) { + // if we are active on the GPU, we need to send any newly initialized inner members to the device + if (m_pointer_record->m_last_space == GPU && old_size < m_elems) { + umpire::ResourceManager & umpire_rm = umpire::ResourceManager::getInstance(); + void *src = (T*)m_pointer_record->m_pointers[CPU] + old_size; + void *dst = (T*)m_pointer_record->m_pointers[GPU] + old_size; + umpire_rm.copy(dst,src,(m_elems-old_size)*sizeof(T)); + } + } CHAI_LOG(Debug, "m_active_ptr reallocated at address: " << m_active_pointer); } diff --git a/src/util/forall.hpp b/src/util/forall.hpp index e992b0c5..3121e792 100644 --- a/src/util/forall.hpp +++ b/src/util/forall.hpp @@ -25,7 +25,7 @@ struct gpu { template void forall_kernel_cpu(int begin, int end, LOOP_BODY body) { - for (int i = 0; i < (end - begin); ++i) { + for (int i = begin; i < end; ++i) { body(i); } } @@ -56,7 +56,7 @@ __global__ void forall_kernel_gpu(int start, int length, LOOP_BODY body) int idx = blockDim.x * blockIdx.x + threadIdx.x; if (idx < length) { - body(idx); + body(idx+start); } } diff --git a/tests/integration/managed_array_tests.cpp b/tests/integration/managed_array_tests.cpp index 8ff1be63..078b0d94 100644 --- a/tests/integration/managed_array_tests.cpp +++ b/tests/integration/managed_array_tests.cpp @@ -1167,6 +1167,65 @@ GPU_TEST(ManagedArray, CallBackConstArrayConst) assert_empty_map(true); } +GPU_TEST(ManagedArray, DeviceInitializedNestedArrays) +{ + int N = 5; + /* Create the outer array. */ + chai::ManagedArray> outerArray(N); + + forall(gpu(), 0, N, + [=]__device__(int i) + { + outerArray[i] = nullptr; + } + ); + + forall(sequential(), 0, N, + [=](int i) + { + outerArray[i] = chai::ManagedArray(1); + } + ); + + forall(gpu(), 0, N, + [=]__device__(int i) + { + for (int j = 0; j < 1; ++j) { + outerArray[i][j] = 0; + } + } + ); + + outerArray.reallocate(2*N); + + forall(sequential(), N,2*N, + [=](int i) + { + outerArray[i] = chai::ManagedArray(1); + } + ); + + forall(gpu(), N, 2*N, + [=]__device__(int i) + { + for (int j = 0; j < 1; ++j) { + outerArray[i][j] = 0; + } + } + ); + + forall(sequential(), 0, 2*N, + [=](int i) + { + for (int j = 0; j < 1; ++j) { + ASSERT_EQ(outerArray[i][j],0); + } + outerArray[i].free(); + } + ); + outerArray.free(); + assert_empty_map(true); +} #endif #endif From b8cf372adf1161ecf86f0506b1ba8a65a4275d0b Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Tue, 28 Apr 2020 15:54:31 -0700 Subject: [PATCH 55/70] adds support for a PINNED Execution space. --- CMakeLists.txt | 3 +- examples/CMakeLists.txt | 5 +++ examples/pinned.cpp | 87 ++++++++++++++++++++++++++++++++++++ src/chai/ArrayManager.cpp | 48 +++++++++++++++++--- src/chai/ArrayManager.hpp | 4 ++ src/chai/CMakeLists.txt | 1 + src/chai/ExecutionSpaces.hpp | 6 +++ src/chai/ManagedArray.hpp | 2 + src/chai/ManagedArray.inl | 52 ++++++++++++--------- src/chai/config.hpp.in | 1 + src/util/forall.hpp | 21 +++++++++ 11 files changed, 203 insertions(+), 27 deletions(-) create mode 100644 examples/pinned.cpp diff --git a/CMakeLists.txt b/CMakeLists.txt index 8fa9af23..bd5985d2 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -12,13 +12,14 @@ project(Chai LANGUAGES CXX VERSION 2.0.0) set(ENABLE_CUDA Off CACHE BOOL "Enable CUDA") set(ENABLE_HIP Off CACHE BOOL "Enable HIP") set(ENABLE_GPU_SIMULATION_MODE Off CACHE BOOL "Enable GPU Simulation Mode") -set(ENABLE_OPENMP On CACHE BOOL "Enable OpenMP") +set(ENABLE_OPENMP OFF CACHE BOOL "Enable OpenMP") set(ENABLE_MPI Off CACHE BOOL "Enable MPI (for umpire replay only)") set(ENABLE_BENCHMARKS On CACHE BOOL "Enable benchmarks") option(ENABLE_IMPLICIT_CONVERSIONS "Enable implicit conversions to-from raw pointers" On) option(DISABLE_RM "Make ManagedArray a thin wrapper" Off) mark_as_advanced(DISABLE_RM) option(ENABLE_UM "Use CUDA unified (managed) memory" Off) +option(ENABLE_PINNED "Use pinned host memory" Off) option(ENABLE_RAJA_PLUGIN "Build plugin to set RAJA execution spaces" Off) option(CHAI_ENABLE_GPU_ERROR_CHECKING "Enable GPU error checking" On) option(CHAI_DEBUG "Enable Debug Logging.") diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt index e380076e..b4c83bf8 100644 --- a/examples/CMakeLists.txt +++ b/examples/CMakeLists.txt @@ -28,5 +28,10 @@ if (ENABLE_CUDA OR ENABLE_HIP) NAME chai-example.exe SOURCES example.cpp DEPENDS_ON ${chai_umpire_example_depends}) + + blt_add_executable( + NAME pinned.exe + SOURCES pinned.cpp + DEPENDS_ON ${chai_umpire_example_depends}) endif () diff --git a/examples/pinned.cpp b/examples/pinned.cpp new file mode 100644 index 00000000..a1e13e62 --- /dev/null +++ b/examples/pinned.cpp @@ -0,0 +1,87 @@ +////////////////////////////////////////////////////////////////////////////// +// Copyright (c) 2016-19, Lawrence Livermore National Security, LLC and CHAI +// project contributors. See the COPYRIGHT file for details. +// +// SPDX-License-Identifier: BSD-3-Clause +////////////////////////////////////////////////////////////////////////////// +#include "../src/util/forall.hpp" +#include "chai/ManagedArray.hpp" + +#include + +int main(int CHAI_UNUSED_ARG(argc), char** CHAI_UNUSED_ARG(argv)) +{ + std::cout << "Creating a pinned array..." << std::endl; + + chai::ManagedArray array(10, chai::PINNED); + + std::cout << "Setting array on host." << std::endl; + std::cout << "array = ["; + forall(sequential(), 0, 10, [=](int i) { + array[i] = static_cast(i * 1.0f); + std::cout << " " << array[i]; + }); + std::cout << " ]" << std::endl; + + + std::cout << "Doubling on device." << std::endl; + forall(gpu_async(), 0, 10, [=] __device__(int i) { array[i] *= 2.0f; }); + + std::cout << "array = ["; + forall(sequential(), 0, 10, [=](int i) { + std::cout << " " << array[i]; + }); + std::cout << " ]" << std::endl; + + array.free(); + + chai::ManagedArray array_one(4096, chai::PINNED); + chai::ManagedArray array_two(4096, chai::PINNED); + chai::ManagedArray array_three(4096, chai::PINNED); + chai::ManagedArray array_four(4096, chai::PINNED); + chai::ManagedArray array_five(4096, chai::PINNED); + + std::cout << "Setting arrays on host." << std::endl; + forall(sequential(), 0, 4096, [=](int i) { + array_one[i] = static_cast(i * 1.0f); + array_two[i] = static_cast(i * 2.0f); + array_three[i] = static_cast(i * 3.0f); + array_four[i] = static_cast(i * 4.0f); + array_five[i] = static_cast(i * 5.0f); + }); + + forall(sequential(), 0, 3, [=](int i) { + std::cout << array_one[i] << " "; + std::cout << array_two[i] << " "; + std::cout << array_three[i] << " "; + std::cout << array_four[i] << " "; + std::cout << array_five[i] << " "; + }); + std::cout << std::endl; + + std::cout << "Doubling on device." << std::endl; + forall(gpu_async(), 0, 4096, [=] __device__(int i) { + array_one[i] *= 2.0f; + array_two[i] *= 2.0f; + array_three[i] *= 2.0f; + array_four[i] *= 2.0f; + array_five[i] *= 2.0f; + }); + + forall(sequential(), 0, 3, [=](int i) { + std::cout << array_one[i] << " "; + std::cout << array_two[i] << " "; + std::cout << array_three[i] << " "; + std::cout << array_four[i] << " "; + std::cout << array_five[i] << " "; + }); + std::cout << std::endl; + + array_one.free(); + array_two.free(); + array_three.free(); + array_four.free(); + array_five.free(); + + return 0; +} diff --git a/src/chai/ArrayManager.cpp b/src/chai/ArrayManager.cpp index eccb0d66..c15aa81e 100644 --- a/src/chai/ArrayManager.cpp +++ b/src/chai/ArrayManager.cpp @@ -8,6 +8,10 @@ #include "chai/config.hpp" +#if defined(CHAI_ENABLE_CUDA) +#include "cuda_runtime_api.h" +#endif + #include "umpire/ResourceManager.hpp" namespace chai @@ -48,6 +52,10 @@ ArrayManager::ArrayManager() : m_allocators[UM] = new umpire::Allocator(m_resource_manager.getAllocator("UM")); #endif +#if defined(CHAI_ENABLE_PINNED) + m_allocators[PINNED] = + new umpire::Allocator(m_resource_manager.getAllocator("PINNED")); +#endif } void ArrayManager::registerPointer( @@ -147,6 +155,12 @@ void ArrayManager::setExecutionSpace(ExecutionSpace space) CHAI_LOG(Debug, "Setting execution space to " << space); std::lock_guard lock(m_mutex); +#if defined(CHAI_ENABLE_PINNED) + if (chai::GPU == space) { + m_need_sync_for_pinned = true; + } +#endif + m_current_execution_space = space; } @@ -209,17 +223,27 @@ void ArrayManager::move(PointerRecord* record, ExecutionSpace space) return; } + callback(record, ACTION_CAPTURED, space); + + if (space == record->m_last_space) { + return; + } + #if defined(CHAI_ENABLE_UM) if (record->m_last_space == UM) { return; } #endif - callback(record, ACTION_CAPTURED, space); - - if (space == record->m_last_space) { +#if defined(CHAI_ENABLE_PINNED) + if (record->m_last_space == PINNED) { + if (space == CPU && m_need_sync_for_pinned) { + m_need_sync_for_pinned = false; + cudaDeviceSynchronize(); + } return; } +#endif void* src_pointer = record->m_pointers[record->m_last_space]; void* dst_pointer = record->m_pointers[space]; @@ -273,8 +297,22 @@ void ArrayManager::free(PointerRecord* pointer_record, ExecutionSpace spaceToFre ACTION_FREE, ExecutionSpace(UM)); + auto alloc = m_resource_manager.getAllocator(pointer_record->m_allocators[UM]); + alloc.deallocate(space_ptr); + + for (int space_t = CPU; space_t < NUM_EXECUTION_SPACES; ++space_t) { + if (space_ptr == pointer_record->m_pointers[space_t]) + pointer_record->m_pointers[space_t] = nullptr; + } + } else { +#elif defined(CHAI_ENABLE_PINNED) + if (space_ptr == pointer_record->m_pointers[PINNED]) { + callback(pointer_record, + ACTION_FREE, + ExecutionSpace(PINNED)); + auto alloc = m_resource_manager.getAllocator( - pointer_record->m_allocators[space]); + pointer_record->m_allocators[PINNED]); alloc.deallocate(space_ptr); for (int space_t = CPU; space_t < NUM_EXECUTION_SPACES; ++space_t) { @@ -292,7 +330,7 @@ void ArrayManager::free(PointerRecord* pointer_record, ExecutionSpace spaceToFre alloc.deallocate(space_ptr); pointer_record->m_pointers[space] = nullptr; -#if defined(CHAI_ENABLE_UM) +#if defined(CHAI_ENABLE_UM) || defined(CHAI_ENABLE_PINNED) } #endif } diff --git a/src/chai/ArrayManager.hpp b/src/chai/ArrayManager.hpp index e576a210..188e7cf3 100644 --- a/src/chai/ArrayManager.hpp +++ b/src/chai/ArrayManager.hpp @@ -422,6 +422,10 @@ class ArrayManager * Whether or not to synchronize on device after every CHAI kernel. */ bool m_device_synchronize = false; + +#if defined(CHAI_ENABLE_PINNED) + bool m_need_sync_for_pinned = false; +#endif }; } // end of namespace chai diff --git a/src/chai/CMakeLists.txt b/src/chai/CMakeLists.txt index 027abf9d..e77f09ab 100644 --- a/src/chai/CMakeLists.txt +++ b/src/chai/CMakeLists.txt @@ -12,6 +12,7 @@ set(CHAI_DISABLE_RM ${DISABLE_RM}) set(CHAI_ENABLE_UM ${ENABLE_UM}) set(CHAI_ENABLE_RAJA_PLUGIN ${ENABLE_RAJA_PLUGIN}) set(CHAI_ENABLE_GPU_SIMULATION_MODE ${ENABLE_GPU_SIMULATION_MODE}) +set(CHAI_ENABLE_PINNED ${ENABLE_PINNED}) configure_file( ${PROJECT_SOURCE_DIR}/src/chai/config.hpp.in diff --git a/src/chai/ExecutionSpaces.hpp b/src/chai/ExecutionSpaces.hpp index 9fa32177..dfc7b773 100644 --- a/src/chai/ExecutionSpaces.hpp +++ b/src/chai/ExecutionSpaces.hpp @@ -26,6 +26,9 @@ enum ExecutionSpace { #endif #if defined(CHAI_ENABLE_UM) UM, +#endif +#if defined(CHAI_ENABLE_PINNED) + PINNED, #endif // NUM_EXECUTION_SPACES should always be last! /*! Used to count total number of spaces */ @@ -36,6 +39,9 @@ enum ExecutionSpace { #if !defined(CHAI_ENABLE_UM) ,UM #endif +#if !defined(CHAI_ENABLE_PINNED) + ,PINNED +#endif }; } // end of namespace chai diff --git a/src/chai/ManagedArray.hpp b/src/chai/ManagedArray.hpp index ef21b3c1..51bb45a6 100644 --- a/src/chai/ManagedArray.hpp +++ b/src/chai/ManagedArray.hpp @@ -367,6 +367,8 @@ class ManagedArray : public CHAICopyable #endif } + CHAI_HOST_DEVICE bool isSlice() { return m_is_slice;} + private: CHAI_HOST void modify(size_t i, const T& val) const; diff --git a/src/chai/ManagedArray.inl b/src/chai/ManagedArray.inl index d9fc3ba2..266336a0 100644 --- a/src/chai/ManagedArray.inl +++ b/src/chai/ManagedArray.inl @@ -57,12 +57,6 @@ CHAI_HOST_DEVICE ManagedArray::ManagedArray( this->allocate(elems, space); -#if defined(CHAI_ENABLE_UM) - if(space == UM) { - m_pointer_record->m_pointers[CPU] = m_active_pointer; - m_pointer_record->m_pointers[GPU] = m_active_pointer; - } -#endif #endif } @@ -77,13 +71,6 @@ CHAI_HOST_DEVICE ManagedArray::ManagedArray( { #if !defined(__CUDA_ARCH__) && !defined(__HIP_DEVICE_COMPILE__) this->allocate(elems, space); - - #if defined(CHAI_ENABLE_UM) - if(space == UM) { - m_pointer_record->m_pointers[CPU] = m_active_base_pointer; - m_pointer_record->m_pointers[GPU] = m_active_base_pointer; - } - #endif #endif } @@ -129,7 +116,7 @@ CHAI_HOST_DEVICE ManagedArray::ManagedArray(ManagedArray const& other): if (m_pointer_record) { m_elems = m_pointer_record->m_size/sizeof(T); } - move(); + move(m_resource_manager->getExecutionSpace()); } #endif } @@ -173,8 +160,9 @@ CHAI_HOST void ManagedArray::allocate( if (m_pointer_record == &ArrayManager::s_null_record) { // since we are about to allocate, this will get registered m_pointer_record = new PointerRecord(); - for (int space = CPU; space < NUM_EXECUTION_SPACES; ++space) { - m_pointer_record->m_allocators[space] = m_resource_manager->getAllocatorId(ExecutionSpace(space)); + for (int s = CPU; s < NUM_EXECUTION_SPACES; ++s) { + ExecutionSpace allocator_space = space == PINNED ? PINNED : ExecutionSpace(s); + m_pointer_record->m_allocators[s] = m_resource_manager->getAllocatorId(allocator_space); } } @@ -191,9 +179,23 @@ CHAI_HOST void ManagedArray::allocate( // ManagedArrays to nullptr at allocation, since it is extremely easy to // trigger a moveInnerImpl, which expects inner values to be initialized. initInner(); - + +#if defined(CHAI_ENABLE_UM) + if(space == UM) { + m_pointer_record->m_last_space = UM; + m_pointer_record->m_pointers[CPU] = m_active_pointer; + m_pointer_record->m_pointers[GPU] = m_active_pointer; + } +#endif +#if defined(CHAI_ENABLE_PINNED) + if (space == PINNED) { + m_pointer_record->m_last_space = PINNED; + m_pointer_record->m_pointers[CPU] = m_active_pointer; + m_pointer_record->m_pointers[GPU] = m_active_pointer; + } +#endif CHAI_LOG(Debug, "m_active_base_ptr allocated at address: " << m_active_base_pointer); - } + } } } @@ -393,16 +395,24 @@ void ManagedArray::move(ExecutionSpace space) const m_active_pointer = m_active_base_pointer + m_offset; CHAI_LOG(Debug, "Moved to " << m_active_pointer); +#if defined(CHAI_ENABLE_UM) + if (m_pointer_record->m_last_space == UM) { + } else +#endif +#if defined(CHAI_ENABLE_PINNED) + if (m_pointer_record->m_last_space == PINNED) { + } else +#endif if (!std::is_const::value) { CHAI_LOG(Debug, "T is non-const, registering touch of pointer" << m_active_pointer); m_resource_manager->registerTouch(m_pointer_record); } - if (prev_space == GPU) { + if (space != GPU && prev_space == GPU) { /// Move nested ManagedArrays after the move, so they are working with a valid m_active_pointer for the host, // and so the meta data associated with them are updated with live GPU data moveInnerImpl(); } - } + } } template @@ -423,7 +433,7 @@ CHAI_HOST_DEVICE ManagedArray::operator T*() const { } ExecutionSpace prev_space = m_resource_manager->getExecutionSpace(); m_resource_manager->setExecutionSpace(CPU); - move(); + move(CPU); // always touch regarless of constness of type (don't trust the application not to const-cast) m_resource_manager->registerTouch(m_pointer_record); diff --git a/src/chai/config.hpp.in b/src/chai/config.hpp.in index 40ba15d5..4da4a8e3 100644 --- a/src/chai/config.hpp.in +++ b/src/chai/config.hpp.in @@ -17,5 +17,6 @@ #cmakedefine CHAI_ENABLE_GPU_ERROR_CHECKING #cmakedefine CHAI_ENABLE_RAJA_PLUGIN #cmakedefine CHAI_ENABLE_GPU_SIMULATION_MODE +#cmakedefine CHAI_ENABLE_PINNED #endif // CHAI_config_HPP diff --git a/src/util/forall.hpp b/src/util/forall.hpp index 3121e792..652d412d 100644 --- a/src/util/forall.hpp +++ b/src/util/forall.hpp @@ -20,6 +20,8 @@ struct sequential { #if defined(CHAI_ENABLE_CUDA) || defined(CHAI_ENABLE_HIP) struct gpu { }; + +struct gpu_async {}; #endif template @@ -60,6 +62,25 @@ __global__ void forall_kernel_gpu(int start, int length, LOOP_BODY body) } } +template +void forall(gpu_async, int begin, int end, LOOP_BODY&& body) +{ + chai::ArrayManager* rm = chai::ArrayManager::getInstance(); + + rm->setExecutionSpace(chai::GPU); + + size_t blockSize = 32; + size_t gridSize = (end - begin + blockSize - 1) / blockSize; + +#if defined(CHAI_ENABLE_CUDA) + forall_kernel_gpu<<>>(begin, end - begin, body); +#elif defined(CHAI_ENABLE_HIP) + hipLaunchKernelGGL(forall_kernel_gpu, dim3(gridSize), dim3(blockSize), 0,0, + begin, end - begin, body); +#endif + rm->setExecutionSpace(chai::NONE); +} + /* * \brief Run forall kernel on GPU. */ From c2d05e9289a87efc7d97a711f6acb363ccd69c49 Mon Sep 17 00:00:00 2001 From: Alan Dayton Date: Wed, 29 Apr 2020 09:54:49 -0700 Subject: [PATCH 56/70] Report the correct space where data is leaked --- src/chai/ArrayManager.cpp | 9 ++++++++- 1 file changed, 8 insertions(+), 1 deletion(-) diff --git a/src/chai/ArrayManager.cpp b/src/chai/ArrayManager.cpp index eccb0d66..29d3ac15 100644 --- a/src/chai/ArrayManager.cpp +++ b/src/chai/ArrayManager.cpp @@ -453,7 +453,14 @@ size_t ArrayManager::getTotalSize() const void ArrayManager::reportLeaks() const { for (auto entry : m_pointer_map) { - callback(*entry.second, ACTION_LEAKED, NONE); + const void* pointer = entry.first; + const PointerRecord* record = *entry.second; + + for (int s = CPU; s < NUM_EXECUTION_SPACES; ++s) { + if (pointer == record->m_pointers[s]) { + callback(record, ACTION_LEAKED, ExecutionSpace(s)); + } + } } } From a3c9ac176b953794542892e53deca8d34f1a1195 Mon Sep 17 00:00:00 2001 From: Alan Dayton Date: Wed, 29 Apr 2020 10:51:01 -0700 Subject: [PATCH 57/70] Use const auto& in range for loops --- src/chai/ArrayManager.cpp | 10 +++++----- src/chai/ManagedArray.inl | 2 +- src/chai/ManagedArray_thin.inl | 2 +- 3 files changed, 7 insertions(+), 7 deletions(-) diff --git a/src/chai/ArrayManager.cpp b/src/chai/ArrayManager.cpp index 29d3ac15..3e30c611 100644 --- a/src/chai/ArrayManager.cpp +++ b/src/chai/ArrayManager.cpp @@ -427,7 +427,7 @@ ArrayManager::getPointerMap() const std::unordered_map mapCopy; std::lock_guard lock(m_mutex); - for (auto entry : m_pointer_map) { + for (const auto& entry : m_pointer_map) { mapCopy[entry.first] = *entry.second; } @@ -443,7 +443,7 @@ size_t ArrayManager::getTotalSize() const size_t total = 0; std::lock_guard lock(m_mutex); - for (auto entry : m_pointer_map) { + for (const auto& entry : m_pointer_map) { total += (*entry.second)->m_size; } @@ -452,7 +452,7 @@ size_t ArrayManager::getTotalSize() const void ArrayManager::reportLeaks() const { - for (auto entry : m_pointer_map) { + for (const auto& entry : m_pointer_map) { const void* pointer = entry.first; const PointerRecord* record = *entry.second; @@ -493,7 +493,7 @@ void ArrayManager::evict(ExecutionSpace space, ExecutionSpace destinationSpace) // Now move and evict std::vector pointersToEvict; - for (auto entry : m_pointer_map) { + for (const auto& entry : m_pointer_map) { // Get the pointer record auto record = *entry.second; @@ -511,7 +511,7 @@ void ArrayManager::evict(ExecutionSpace space, ExecutionSpace destinationSpace) // This must be done in a second pass because free erases from m_pointer_map, // which would invalidate the iterator in the above loop - for (auto entry : pointersToEvict) { + for (const auto& entry : pointersToEvict) { free(entry, space); } } diff --git a/src/chai/ManagedArray.inl b/src/chai/ManagedArray.inl index d9fc3ba2..70cc21df 100644 --- a/src/chai/ManagedArray.inl +++ b/src/chai/ManagedArray.inl @@ -39,7 +39,7 @@ CHAI_HOST_DEVICE ManagedArray::ManagedArray( #if !defined(__CUDA_ARCH__) && !defined(__HIP_DEVICE_COMPILE__) m_pointer_record = new PointerRecord(); int i = 0; - for (auto& space : spaces) { + for (const auto& space : spaces) { m_pointer_record->m_allocators[space] = allocators.begin()[i++].getId(); } #endif diff --git a/src/chai/ManagedArray_thin.inl b/src/chai/ManagedArray_thin.inl index 77bb0d9e..07e07f2d 100644 --- a/src/chai/ManagedArray_thin.inl +++ b/src/chai/ManagedArray_thin.inl @@ -24,7 +24,7 @@ CHAI_INLINE CHAI_HOST_DEVICE ManagedArray::ManagedArray( if (m_pointer_record) { int i = 0; - for (auto& space : spaces) { + for (const auto& space : spaces) { m_pointer_record->m_allocators[space] = allocators.begin()[i++].getId(); } } From 55097c99da09c274774fac5305807398c4ce2eb7 Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Wed, 29 Apr 2020 13:32:08 -0700 Subject: [PATCH 58/70] add GPU_SIM_MODE_CHECK. --- src/util/forall.hpp | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/src/util/forall.hpp b/src/util/forall.hpp index 652d412d..ad0ce770 100644 --- a/src/util/forall.hpp +++ b/src/util/forall.hpp @@ -71,8 +71,9 @@ void forall(gpu_async, int begin, int end, LOOP_BODY&& body) size_t blockSize = 32; size_t gridSize = (end - begin + blockSize - 1) / blockSize; - -#if defined(CHAI_ENABLE_CUDA) +#if defined(CHAI_ENABLE_GPU_SIMULATION_MODE) + forall_kernel_cpu(begin, end, body); +#elif defined(CHAI_ENABLE_CUDA) forall_kernel_gpu<<>>(begin, end - begin, body); #elif defined(CHAI_ENABLE_HIP) hipLaunchKernelGGL(forall_kernel_gpu, dim3(gridSize), dim3(blockSize), 0,0, From 954d75a77a5eda8d19b023cf612be9f2610e72be Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Fri, 1 May 2020 08:22:29 -0700 Subject: [PATCH 59/70] make a chai::synchronize for device synchronization, make raja-chai-nested-tests optional. --- CMakeLists.txt | 1 + src/chai/ArrayManager.cpp | 2 +- src/chai/ArrayManager.hpp | 41 +++++++++++++++++++++++++++++++ src/chai/ManagedArray.inl | 6 ++--- src/chai/ManagedArray_thin.inl | 16 ++++++------ src/chai/managed_ptr.hpp | 42 +++++++++----------------------- tests/integration/CMakeLists.txt | 22 +++++++++-------- 7 files changed, 77 insertions(+), 53 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index bd5985d2..1ffd6a73 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -23,6 +23,7 @@ option(ENABLE_PINNED "Use pinned host memory" Off) option(ENABLE_RAJA_PLUGIN "Build plugin to set RAJA execution spaces" Off) option(CHAI_ENABLE_GPU_ERROR_CHECKING "Enable GPU error checking" On) option(CHAI_DEBUG "Enable Debug Logging.") +set(ENABLE_RAJA_NESTED_TEST ON CACHE BOOL "Enable raja-chai-nested-tests, which fails to build on Debug CUDA builds.") set(ENABLE_TESTS On CACHE BOOL "") set(ENABLE_EXAMPLES On CACHE BOOL "") diff --git a/src/chai/ArrayManager.cpp b/src/chai/ArrayManager.cpp index 03aedc1d..b3fb0f8a 100644 --- a/src/chai/ArrayManager.cpp +++ b/src/chai/ArrayManager.cpp @@ -239,7 +239,7 @@ void ArrayManager::move(PointerRecord* record, ExecutionSpace space) if (record->m_last_space == PINNED) { if (space == CPU && m_need_sync_for_pinned) { m_need_sync_for_pinned = false; - cudaDeviceSynchronize(); + synchronize(); } return; } diff --git a/src/chai/ArrayManager.hpp b/src/chai/ArrayManager.hpp index 188e7cf3..fabb0cec 100644 --- a/src/chai/ArrayManager.hpp +++ b/src/chai/ArrayManager.hpp @@ -24,7 +24,48 @@ namespace chai { +// CHAI_GPU_ERROR_CHECK macro +#if defined(__CUDACC__) || defined(__HIPCC__) +#ifdef CHAI_ENABLE_GPU_ERROR_CHECKING + +#ifdef __CUDACC__ +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 __HIPCC__ +inline void gpuErrorCheck(hipError_t code, const char *file, int line, bool abort=true) +{ + if (code != cudaSuccess) { + 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 + +inline void synchronize() { +#if defined(__HIPCC__) && defined (CHAI_ENABLE_HIP) &&!defined(__HIP_DEVICE_COMPILE__) + CHAI_GPU_ERROR_CHECK(hipDeviceSynchronize()); +#elif defined(__CUDACC__) && defined (CHAI_ENABLE_CUDA) &&!defined(__CUDA_ARCH__) + CHAI_GPU_ERROR_CHECK(cudaDeviceSynchronize()); +#endif +} /*! * \brief Singleton that manages caching and movement of ManagedArray objects. * diff --git a/src/chai/ManagedArray.inl b/src/chai/ManagedArray.inl index 266336a0..365b3446 100644 --- a/src/chai/ManagedArray.inl +++ b/src/chai/ManagedArray.inl @@ -302,7 +302,7 @@ typename ManagedArray::T_non_const ManagedArray::pick(size_t i) const { #if !defined(__CUDA_ARCH__) && !defined(__HIP_DEVICE_COMPILE__) #if defined(CHAI_ENABLE_UM) if(m_pointer_record->m_pointers[UM] == m_active_base_pointer) { - cudaDeviceSynchronize(); + synchronize(); return (T_non_const)(m_active_pointer[i]); } #endif @@ -325,7 +325,7 @@ CHAI_HOST_DEVICE void ManagedArray::set(size_t i, T val) const { #if !defined(__CUDA_ARCH__) && !defined(__HIP_DEVICE_COMPILE__) #if defined(CHAI_ENABLE_UM) if(m_pointer_record->m_pointers[UM] == m_active_pointer) { - cudaDeviceSynchronize(); + synchronize(); m_active_pointer[i] = val; return; } @@ -347,7 +347,7 @@ CHAI_INLINE CHAI_HOST void ManagedArray::modify(size_t i, const T& val) const { #if defined(CHAI_ENABLE_UM) if(m_pointer_record->m_pointers[UM] == m_active_pointer) { - cudaDeviceSynchronize(); + synchronize(); m_active_pointer[i] = m_active_pointer[i] + val; return; } diff --git a/src/chai/ManagedArray_thin.inl b/src/chai/ManagedArray_thin.inl index 77bb0d9e..f4660877 100644 --- a/src/chai/ManagedArray_thin.inl +++ b/src/chai/ManagedArray_thin.inl @@ -193,8 +193,8 @@ template CHAI_INLINE CHAI_HOST_DEVICE typename ManagedArray::T_non_const ManagedArray< T>::pick(size_t i) const { -#if !defined(__CUDA_ARCH__) && defined(CHAI_ENABLE_UM) - cudaDeviceSynchronize(); +#ifdef CHAI_ENABLE_UM + synchronize(); #endif return (T_non_const)m_active_pointer[i]; } @@ -202,8 +202,8 @@ CHAI_INLINE CHAI_HOST_DEVICE typename ManagedArray::T_non_const ManagedArray< template CHAI_INLINE CHAI_HOST_DEVICE void ManagedArray::set(size_t i, T val) const { -#if !defined(__CUDA_ARCH__) && defined(CHAI_ENABLE_UM) - cudaDeviceSynchronize(); +#if defined(CHAI_ENABLE_UM) + synchronize(); #endif m_active_pointer[i] = val; } @@ -211,8 +211,8 @@ CHAI_INLINE CHAI_HOST_DEVICE void ManagedArray::set(size_t i, T val) const template CHAI_INLINE CHAI_HOST_DEVICE void ManagedArray::incr(size_t i) const { -#if !defined(__CUDA_ARCH__) && defined(CHAI_ENABLE_UM) - cudaDeviceSynchronize(); +#if defined(CHAI_ENABLE_UM) + synchronize(); #endif ++m_active_pointer[i]; } @@ -220,8 +220,8 @@ CHAI_INLINE CHAI_HOST_DEVICE void ManagedArray::incr(size_t i) const template CHAI_INLINE CHAI_HOST_DEVICE void ManagedArray::decr(size_t i) const { -#if !defined(__CUDA_ARCH__) && defined(CHAI_ENABLE_UM) - cudaDeviceSynchronize(); +#if defined(CHAI_ENABLE_UM) + synchronize(); #endif --m_active_pointer[i]; } diff --git a/src/chai/managed_ptr.hpp b/src/chai/managed_ptr.hpp index 39dca4b8..3d2a9b2a 100644 --- a/src/chai/managed_ptr.hpp +++ b/src/chai/managed_ptr.hpp @@ -59,26 +59,6 @@ #include #include -#ifdef __CUDACC__ - -#ifdef CHAI_ENABLE_GPU_ERROR_CHECKING - -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); - } - } -} - -#define GPU_ERROR_CHECK(code) { gpuErrorCheck((code), __FILE__, __LINE__); } -#else // CHAI_ENABLE_GPU_ERROR_CHECKING -#define GPU_ERROR_CHECK(code) code -#endif // CHAI_ENABLE_GPU_ERROR_CHECKING - -#endif // __CUDACC__ namespace chai { namespace detail { @@ -153,7 +133,7 @@ namespace chai { /// only be valid in the correct context. Take care when passing raw pointers /// as arguments to member functions. /// Be aware that CHAI checks every CUDA API call for GPU errors by default. To - /// turn off GPU error checking, pass -DCHAI_ENABLE_GPU_ERROR_CHECKING=OFF as + /// turn off GPU error checking, pass -DCHAI_ENABLE_CHAI_GPU_ERROR_CHECKING=OFF as /// an argument to cmake when building CHAI. To turn on synchronization after /// every kernel, call ArrayManager::getInstance()->enableDeviceSynchronize(). /// Alternatively, call cudaDeviceSynchronize() after any call to make_managed, @@ -557,7 +537,7 @@ namespace chai { #ifndef CHAI_DISABLE_RM if (ArrayManager::getInstance()->deviceSynchronize()) { - GPU_ERROR_CHECK(cudaDeviceSynchronize()); + synchronize(); } #endif } @@ -589,7 +569,7 @@ namespace chai { #ifndef CHAI_DISABLE_RM if (ArrayManager::getInstance()->deviceSynchronize()) { - GPU_ERROR_CHECK(cudaDeviceSynchronize()); + synchronize(); } #endif } @@ -929,20 +909,20 @@ namespace chai { // Allocate space on the GPU to hold the pointer to the new object T** gpuBuffer; - GPU_ERROR_CHECK(cudaMalloc(&gpuBuffer, sizeof(T*))); + CHAI_GPU_ERROR_CHECK(cudaMalloc(&gpuBuffer, sizeof(T*))); // Create the object on the device make_on_device<<<1, 1>>>(gpuBuffer, args...); #ifndef CHAI_DISABLE_RM if (ArrayManager::getInstance()->deviceSynchronize()) { - GPU_ERROR_CHECK(cudaDeviceSynchronize()); + synchronize(); } #endif // Allocate space on the CPU for the pointer and copy the pointer to the CPU T** cpuBuffer = (T**) malloc(sizeof(T*)); - GPU_ERROR_CHECK(cudaMemcpy(cpuBuffer, gpuBuffer, sizeof(T*), + CHAI_GPU_ERROR_CHECK(cudaMemcpy(cpuBuffer, gpuBuffer, sizeof(T*), cudaMemcpyDeviceToHost)); // Get the GPU pointer @@ -950,7 +930,7 @@ namespace chai { // Free the host and device buffers free(cpuBuffer); - GPU_ERROR_CHECK(cudaFree(gpuBuffer)); + CHAI_GPU_ERROR_CHECK(cudaFree(gpuBuffer)); #ifndef CHAI_DISABLE_RM // Set the execution space back to the previous value @@ -987,20 +967,20 @@ namespace chai { // Allocate space on the GPU to hold the pointer to the new object T** gpuBuffer; - GPU_ERROR_CHECK(cudaMalloc(&gpuBuffer, sizeof(T*))); + CHAI_GPU_ERROR_CHECK(cudaMalloc(&gpuBuffer, sizeof(T*))); // Create the object on the device make_on_device_from_factory<<<1, 1>>>(gpuBuffer, f, args...); #ifndef CHAI_DISABLE_RM if (ArrayManager::getInstance()->deviceSynchronize()) { - GPU_ERROR_CHECK(cudaDeviceSynchronize()); + synchronize(); } #endif // Allocate space on the CPU for the pointer and copy the pointer to the CPU T** cpuBuffer = (T**) malloc(sizeof(T*)); - GPU_ERROR_CHECK(cudaMemcpy(cpuBuffer, gpuBuffer, sizeof(T*), + CHAI_GPU_ERROR_CHECK(cudaMemcpy(cpuBuffer, gpuBuffer, sizeof(T*), cudaMemcpyDeviceToHost)); // Get the GPU pointer @@ -1008,7 +988,7 @@ namespace chai { // Free the host and device buffers free(cpuBuffer); - GPU_ERROR_CHECK(cudaFree(gpuBuffer)); + CHAI_GPU_ERROR_CHECK(cudaFree(gpuBuffer)); #ifndef CHAI_DISABLE_RM // Set the execution space back to the previous value diff --git a/tests/integration/CMakeLists.txt b/tests/integration/CMakeLists.txt index 3dcc2c31..21e2a1e3 100644 --- a/tests/integration/CMakeLists.txt +++ b/tests/integration/CMakeLists.txt @@ -56,16 +56,18 @@ if (ENABLE_RAJA_PLUGIN) raja-chai-tests PUBLIC ${PROJECT_BINARY_DIR}/include) - blt_add_executable( - NAME raja-chai-nested-tests - SOURCES raja-chai-nested.cpp - DEPENDS_ON ${raja_test_depends}) + if (ENABLE_RAJA_NESTED_TEST) + blt_add_executable( + NAME raja-chai-nested-tests + SOURCES raja-chai-nested.cpp + DEPENDS_ON ${raja_test_depends}) - blt_add_test( - NAME raja-chai-nested-tests - COMMAND raja-chai-nested-tests) + blt_add_test( + NAME raja-chai-nested-tests + COMMAND raja-chai-nested-tests) - target_include_directories( - raja-chai-nested-tests - PUBLIC ${PROJECT_BINARY_DIR}/include) + target_include_directories( + raja-chai-nested-tests + PUBLIC ${PROJECT_BINARY_DIR}/include) + endif () endif () From 49843eb1346d04fb719daff501b7cc2f1adcedcf Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Fri, 1 May 2020 13:22:07 -0700 Subject: [PATCH 60/70] add some documentation. --- src/chai/ArrayManager.hpp | 3 +++ src/chai/ManagedArray.hpp | 3 +++ 2 files changed, 6 insertions(+) diff --git a/src/chai/ArrayManager.hpp b/src/chai/ArrayManager.hpp index fabb0cec..9534894e 100644 --- a/src/chai/ArrayManager.hpp +++ b/src/chai/ArrayManager.hpp @@ -465,6 +465,9 @@ class ArrayManager bool m_device_synchronize = false; #if defined(CHAI_ENABLE_PINNED) + /*! + * Whether or not a synchronize is needed to ensure pinned memory is up to date. + */ bool m_need_sync_for_pinned = false; #endif }; diff --git a/src/chai/ManagedArray.hpp b/src/chai/ManagedArray.hpp index 51bb45a6..5af7cec1 100644 --- a/src/chai/ManagedArray.hpp +++ b/src/chai/ManagedArray.hpp @@ -367,6 +367,9 @@ class ManagedArray : public CHAICopyable #endif } + /*! + * Accessor for m_is_slice -whether this array was created with a slice() command. + */ CHAI_HOST_DEVICE bool isSlice() { return m_is_slice;} From e37aca33fd58ae827c8ac92c29c72a833646ea5b Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Sat, 2 May 2020 14:25:00 -0700 Subject: [PATCH 61/70] ArrayManager is not actually built by nvcc, so we were never synchronizing for calls to chai::synchronize from the Array Manager. --- src/chai/ArrayManager.hpp | 16 +++++++++++----- 1 file changed, 11 insertions(+), 5 deletions(-) diff --git a/src/chai/ArrayManager.hpp b/src/chai/ArrayManager.hpp index 9534894e..9353fe1f 100644 --- a/src/chai/ArrayManager.hpp +++ b/src/chai/ArrayManager.hpp @@ -22,14 +22,20 @@ #include "umpire/Allocator.hpp" #include "umpire/util/MemoryMap.hpp" +#if defined(CHAI_ENABLE_CUDA) +#include +#endif +#if defined(CHAI_ENABLE_HIP) +#include +#endif namespace chai { // CHAI_GPU_ERROR_CHECK macro -#if defined(__CUDACC__) || defined(__HIPCC__) +#if defined(CHAI_ENABLE_CUDA) || defined(CHAI_ENABLE_HIP) #ifdef CHAI_ENABLE_GPU_ERROR_CHECKING -#ifdef __CUDACC__ +#ifdef CHAI_ENABLE_CUDA inline void gpuErrorCheck(cudaError_t code, const char *file, int line, bool abort=true) { if (code != cudaSuccess) { @@ -39,7 +45,7 @@ inline void gpuErrorCheck(cudaError_t code, const char *file, int line, bool abo } } } -#elif defined __HIPCC__ +#elif CHAI_ENABLE_HIP inline void gpuErrorCheck(hipError_t code, const char *file, int line, bool abort=true) { if (code != cudaSuccess) { @@ -60,9 +66,9 @@ inline void gpuErrorCheck(hipError_t code, const char *file, int line, bool abor #endif inline void synchronize() { -#if defined(__HIPCC__) && defined (CHAI_ENABLE_HIP) &&!defined(__HIP_DEVICE_COMPILE__) +#if defined (CHAI_ENABLE_HIP) &&!defined(__HIP_DEVICE_COMPILE__) CHAI_GPU_ERROR_CHECK(hipDeviceSynchronize()); -#elif defined(__CUDACC__) && defined (CHAI_ENABLE_CUDA) &&!defined(__CUDA_ARCH__) +#elif defined (CHAI_ENABLE_CUDA) &&!defined(__CUDA_ARCH__) CHAI_GPU_ERROR_CHECK(cudaDeviceSynchronize()); #endif } From 9259c579a969f0c0fd6c844acff78a10f2f2ee09 Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Wed, 6 May 2020 07:39:21 -0700 Subject: [PATCH 62/70] remove unnecessary moves for pinned memory. --- src/chai/ArrayManager.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/src/chai/ArrayManager.cpp b/src/chai/ArrayManager.cpp index 224cb7a6..c9436bd1 100644 --- a/src/chai/ArrayManager.cpp +++ b/src/chai/ArrayManager.cpp @@ -255,7 +255,8 @@ void ArrayManager::move(PointerRecord* record, ExecutionSpace space) if (!record->m_touched[record->m_last_space]) { return; - } else { + } else if (dst_pointer != src_pointer) { + // Exclude the copy if src and dst are the same (can happen for PINNED memory) { std::lock_guard lock(m_mutex); m_resource_manager.copy(dst_pointer, src_pointer); From 211029d3dd91780f9db73d1215dc4ed84a02d101 Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Wed, 6 May 2020 07:42:21 -0700 Subject: [PATCH 63/70] fix bad CHAI_GPU_ERROR_CHECKING name./ --- src/chai/managed_ptr.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/chai/managed_ptr.hpp b/src/chai/managed_ptr.hpp index 3d2a9b2a..1ec11966 100644 --- a/src/chai/managed_ptr.hpp +++ b/src/chai/managed_ptr.hpp @@ -133,7 +133,7 @@ namespace chai { /// only be valid in the correct context. Take care when passing raw pointers /// as arguments to member functions. /// Be aware that CHAI checks every CUDA API call for GPU errors by default. To - /// turn off GPU error checking, pass -DCHAI_ENABLE_CHAI_GPU_ERROR_CHECKING=OFF as + /// turn off GPU error checking, pass -DCHAI_ENABLE_GPU_ERROR_CHECKING=OFF as /// an argument to cmake when building CHAI. To turn on synchronization after /// every kernel, call ArrayManager::getInstance()->enableDeviceSynchronize(). /// Alternatively, call cudaDeviceSynchronize() after any call to make_managed, From 5ccd026c1cf50e292c13858786d283898aa5b511 Mon Sep 17 00:00:00 2001 From: Alan Dayton Date: Fri, 8 May 2020 09:37:36 -0700 Subject: [PATCH 64/70] Fix GPU simulation mode --- src/chai/ArrayManager.cpp | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) diff --git a/src/chai/ArrayManager.cpp b/src/chai/ArrayManager.cpp index c9436bd1..0884e27d 100644 --- a/src/chai/ArrayManager.cpp +++ b/src/chai/ArrayManager.cpp @@ -53,9 +53,14 @@ ArrayManager::ArrayManager() : new umpire::Allocator(m_resource_manager.getAllocator("UM")); #endif #if defined(CHAI_ENABLE_PINNED) - m_allocators[PINNED] = +#if defined(CHAI_ENABLE_GPU_SIMULATION_MODE) + m_allocators[PINNED] = + new umpire::Allocator(m_resource_manager.getAllocator("HOST")); +#else + m_allocators[PINNED] = new umpire::Allocator(m_resource_manager.getAllocator("PINNED")); #endif +#endif } void ArrayManager::registerPointer( From 8fd29b2b4210300e8344aadc1987b7750219f47e Mon Sep 17 00:00:00 2001 From: Alan Dayton Date: Fri, 8 May 2020 11:39:44 -0700 Subject: [PATCH 65/70] Fix for enabling pinned memory in a host build --- src/chai/ArrayManager.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/src/chai/ArrayManager.cpp b/src/chai/ArrayManager.cpp index 0884e27d..80e5be84 100644 --- a/src/chai/ArrayManager.cpp +++ b/src/chai/ArrayManager.cpp @@ -52,8 +52,9 @@ ArrayManager::ArrayManager() : m_allocators[UM] = new umpire::Allocator(m_resource_manager.getAllocator("UM")); #endif + #if defined(CHAI_ENABLE_PINNED) -#if defined(CHAI_ENABLE_GPU_SIMULATION_MODE) +#if (defined(CHAI_ENABLE_CUDA) || defined(CHAI_ENABLE_HIP)) && !defined(CHAI_ENABLE_GPU_SIMULATION_MODE) m_allocators[PINNED] = new umpire::Allocator(m_resource_manager.getAllocator("HOST")); #else From e9dcf1655263c62992bd2de69dc4f36e7707341d Mon Sep 17 00:00:00 2001 From: Alan Dayton Date: Fri, 8 May 2020 16:03:55 -0700 Subject: [PATCH 66/70] Fix host only build when pinned memory is enabled --- src/chai/ManagedArray.inl | 3 +++ 1 file changed, 3 insertions(+) diff --git a/src/chai/ManagedArray.inl b/src/chai/ManagedArray.inl index 75b1a01d..42081511 100644 --- a/src/chai/ManagedArray.inl +++ b/src/chai/ManagedArray.inl @@ -191,7 +191,10 @@ CHAI_HOST void ManagedArray::allocate( if (space == PINNED) { m_pointer_record->m_last_space = PINNED; m_pointer_record->m_pointers[CPU] = m_active_pointer; + +#if defined(CHAI_ENABLE_CUDA) || defined(CHAI_ENABLE_HIP) || defined(CHAI_ENABLE_GPU_SIMULATION_MODE) m_pointer_record->m_pointers[GPU] = m_active_pointer; +#endif } #endif CHAI_LOG(Debug, "m_active_base_ptr allocated at address: " << m_active_base_pointer); From 4ce5d6d4bd727c6411db3833d02595d525a72e48 Mon Sep 17 00:00:00 2001 From: Alan Dayton Date: Fri, 8 May 2020 16:05:37 -0700 Subject: [PATCH 67/70] Fix host only build when um memory is enabled --- src/chai/ManagedArray.inl | 3 +++ 1 file changed, 3 insertions(+) diff --git a/src/chai/ManagedArray.inl b/src/chai/ManagedArray.inl index 42081511..09e84a52 100644 --- a/src/chai/ManagedArray.inl +++ b/src/chai/ManagedArray.inl @@ -184,7 +184,10 @@ CHAI_HOST void ManagedArray::allocate( if(space == UM) { m_pointer_record->m_last_space = UM; m_pointer_record->m_pointers[CPU] = m_active_pointer; + +#if defined(CHAI_ENABLE_CUDA) || defined(CHAI_ENABLE_HIP) || defined(CHAI_ENABLE_GPU_SIMULATION_MODE) m_pointer_record->m_pointers[GPU] = m_active_pointer; +#endif } #endif #if defined(CHAI_ENABLE_PINNED) From c5fe7dcfba7dc83a6e8eb06b59e6bf6db83a7de5 Mon Sep 17 00:00:00 2001 From: David Beckingsale Date: Wed, 13 May 2020 14:10:00 -0700 Subject: [PATCH 68/70] Bump umpire and scripts --- scripts/make_release_tarball.sh | 2 +- src/tpl/umpire | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/scripts/make_release_tarball.sh b/scripts/make_release_tarball.sh index 6a062604..6258c5b8 100755 --- a/scripts/make_release_tarball.sh +++ b/scripts/make_release_tarball.sh @@ -7,7 +7,7 @@ ############################################################################## TAR_CMD=gtar -VERSION=1.2.0 +VERSION=2.1.0 git archive --prefix=chai-${VERSION}/ -o chai-${VERSION}.tar HEAD 2> /dev/null diff --git a/src/tpl/umpire b/src/tpl/umpire index 3db26e6a..52e10c05 160000 --- a/src/tpl/umpire +++ b/src/tpl/umpire @@ -1 +1 @@ -Subproject commit 3db26e6a2626ee8c0cfa5c9769cfac6e33587122 +Subproject commit 52e10c05cd40dfdfde186c1e63213695f5aeaf65 From 3b5fb91ca38774a7e311f9b209e7734c00a541cb Mon Sep 17 00:00:00 2001 From: David Beckingsale Date: Wed, 13 May 2020 14:11:26 -0700 Subject: [PATCH 69/70] Bumping version numbers --- CMakeLists.txt | 2 +- README.md | 2 +- docs/sphinx/conf.py | 4 ++-- docs/sphinx/conf.py.in | 4 ++-- 4 files changed, 6 insertions(+), 6 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 1ffd6a73..aa3ad276 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -7,7 +7,7 @@ cmake_policy(SET CMP0057 NEW) cmake_policy(SET CMP0048 NEW) -project(Chai LANGUAGES CXX VERSION 2.0.0) +project(Chai LANGUAGES CXX VERSION 2.1.0) set(ENABLE_CUDA Off CACHE BOOL "Enable CUDA") set(ENABLE_HIP Off CACHE BOOL "Enable HIP") diff --git a/README.md b/README.md index 93a62f23..a616b9c9 100644 --- a/README.md +++ b/README.md @@ -1,4 +1,4 @@ -# CHAI v2.0 +# CHAI v2.1 [![Azure Build Status](https://dev.azure.com/davidbeckingsale/CHAI/_apis/build/status/LLNL.CHAI?branchName=develop)](https://dev.azure.com/davidbeckingsale/CHAI/_build/latest?definitionId=2&branchName=develop) [![Build Status](https://travis-ci.org/LLNL/CHAI.svg?branch=develop)](https://travis-ci.org/LLNL/CHAI) diff --git a/docs/sphinx/conf.py b/docs/sphinx/conf.py index cdf48d8b..404b4881 100644 --- a/docs/sphinx/conf.py +++ b/docs/sphinx/conf.py @@ -61,9 +61,9 @@ # built documents. # # The short X.Y version. -version = u'2.0' +version = u'2.1' # The full version, including alpha/beta/rc tags. -release = u'2.0.0' +release = u'2.1.0' # The language for content autogenerated by Sphinx. Refer to documentation # for a list of supported languages. diff --git a/docs/sphinx/conf.py.in b/docs/sphinx/conf.py.in index 8e1585b0..de241adb 100644 --- a/docs/sphinx/conf.py.in +++ b/docs/sphinx/conf.py.in @@ -60,9 +60,9 @@ author = u'' # built documents. # # The short X.Y version. -version = u'2.0' +version = u'2.1' # The full version, including alpha/beta/rc tags. -release = u'2.0.0' +release = u'2.1.0' # The language for content autogenerated by Sphinx. Refer to documentation # for a list of supported languages. From b7cd2b45ca730194bfbf98ae4e69e22481d8c512 Mon Sep 17 00:00:00 2001 From: David Beckingsale Date: Wed, 13 May 2020 16:18:30 -0700 Subject: [PATCH 70/70] Update copyright year --- CMakeLists.txt | 2 +- benchmarks/CMakeLists.txt | 2 +- benchmarks/chai_arraymanager_benchmarks.cpp | 2 +- benchmarks/chai_benchmark_utils.hpp | 2 +- benchmarks/chai_managedarray_benchmarks.cpp | 2 +- cmake/ChaiBasics.cmake | 2 +- cmake/thirdparty/SetupChaiThirdparty.cmake | 2 +- docs/CMakeLists.txt | 2 +- docs/doxygen/CMakeLists.txt | 2 +- docs/doxygen/Doxyfile.in | 2 +- docs/sphinx/CMakeLists.txt | 2 +- docs/sphinx/conf.py | 2 +- docs/sphinx/conf.py.in | 2 +- examples/CMakeLists.txt | 2 +- examples/chai-umpire-allocators.cpp | 2 +- examples/ex1.cpp | 2 +- examples/example.cpp | 2 +- examples/pinned.cpp | 2 +- scripts/apply-license-info.sh | 2 +- scripts/format-source.sh | 2 +- scripts/license.txt | 2 +- scripts/make_release_tarball.sh | 2 +- scripts/travis/build_and_test.sh | 2 +- scripts/travis/install_llvm.sh | 2 +- scripts/update-copyright-year.sh | 4 ++-- src/CMakeLists.txt | 2 +- src/chai/ArrayManager.cpp | 2 +- src/chai/ArrayManager.hpp | 2 +- src/chai/ArrayManager.inl | 2 +- src/chai/CMakeLists.txt | 2 +- src/chai/ChaiMacros.hpp | 2 +- src/chai/ExecutionSpaces.hpp | 2 +- src/chai/ManagedArray.hpp | 2 +- src/chai/ManagedArray.inl | 2 +- src/chai/ManagedArray_thin.inl | 2 +- src/chai/PointerRecord.hpp | 2 +- src/chai/Types.hpp | 2 +- src/chai/chai-config.cmake.in | 2 +- src/chai/config.hpp.in | 2 +- src/util/forall.hpp | 2 +- tests/CMakeLists.txt | 2 +- tests/integration/CMakeLists.txt | 2 +- tests/integration/managed_array_tests.cpp | 2 +- tests/integration/raja-chai-nested.cpp | 2 +- tests/integration/raja-chai-tests.cpp | 2 +- tests/unit/CMakeLists.txt | 2 +- tests/unit/array_manager_unit_tests.cpp | 2 +- tests/unit/managed_array_unit_tests.cpp | 2 +- 48 files changed, 49 insertions(+), 49 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index aa3ad276..17baa54e 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -1,5 +1,5 @@ ############################################################################## -# Copyright (c) 2016-19, Lawrence Livermore National Security, LLC and CHAI +# Copyright (c) 2016-20, Lawrence Livermore National Security, LLC and CHAI # project contributors. See the COPYRIGHT file for details. # # SPDX-License-Identifier: BSD-3-Clause diff --git a/benchmarks/CMakeLists.txt b/benchmarks/CMakeLists.txt index 0bb27260..14bbc1be 100644 --- a/benchmarks/CMakeLists.txt +++ b/benchmarks/CMakeLists.txt @@ -1,5 +1,5 @@ ############################################################################## -# Copyright (c) 2016-19, Lawrence Livermore National Security, LLC and CHAI +# Copyright (c) 2016-20, Lawrence Livermore National Security, LLC and CHAI # project contributors. See the COPYRIGHT file for details. # # SPDX-License-Identifier: BSD-3-Clause diff --git a/benchmarks/chai_arraymanager_benchmarks.cpp b/benchmarks/chai_arraymanager_benchmarks.cpp index c58a4987..14e9406d 100644 --- a/benchmarks/chai_arraymanager_benchmarks.cpp +++ b/benchmarks/chai_arraymanager_benchmarks.cpp @@ -1,5 +1,5 @@ ////////////////////////////////////////////////////////////////////////////// -// Copyright (c) 2016-19, Lawrence Livermore National Security, LLC and CHAI +// Copyright (c) 2016-20, Lawrence Livermore National Security, LLC and CHAI // project contributors. See the COPYRIGHT file for details. // // SPDX-License-Identifier: BSD-3-Clause diff --git a/benchmarks/chai_benchmark_utils.hpp b/benchmarks/chai_benchmark_utils.hpp index 977f5b65..524f1309 100644 --- a/benchmarks/chai_benchmark_utils.hpp +++ b/benchmarks/chai_benchmark_utils.hpp @@ -1,5 +1,5 @@ ////////////////////////////////////////////////////////////////////////////// -// Copyright (c) 2016-19, Lawrence Livermore National Security, LLC and CHAI +// Copyright (c) 2016-20, Lawrence Livermore National Security, LLC and CHAI // project contributors. See the COPYRIGHT file for details. // // SPDX-License-Identifier: BSD-3-Clause diff --git a/benchmarks/chai_managedarray_benchmarks.cpp b/benchmarks/chai_managedarray_benchmarks.cpp index 4fcb33bf..8e2b6adb 100644 --- a/benchmarks/chai_managedarray_benchmarks.cpp +++ b/benchmarks/chai_managedarray_benchmarks.cpp @@ -1,5 +1,5 @@ ////////////////////////////////////////////////////////////////////////////// -// Copyright (c) 2016-19, Lawrence Livermore National Security, LLC and CHAI +// Copyright (c) 2016-20, Lawrence Livermore National Security, LLC and CHAI // project contributors. See the COPYRIGHT file for details. // // SPDX-License-Identifier: BSD-3-Clause diff --git a/cmake/ChaiBasics.cmake b/cmake/ChaiBasics.cmake index 93db482b..a2a96026 100644 --- a/cmake/ChaiBasics.cmake +++ b/cmake/ChaiBasics.cmake @@ -1,5 +1,5 @@ ############################################################################## -# Copyright (c) 2016-19, Lawrence Livermore National Security, LLC and CHAI +# Copyright (c) 2016-20, Lawrence Livermore National Security, LLC and CHAI # project contributors. See the COPYRIGHT file for details. # # SPDX-License-Identifier: BSD-3-Clause diff --git a/cmake/thirdparty/SetupChaiThirdparty.cmake b/cmake/thirdparty/SetupChaiThirdparty.cmake index 492a62d9..6b6b6ebf 100644 --- a/cmake/thirdparty/SetupChaiThirdparty.cmake +++ b/cmake/thirdparty/SetupChaiThirdparty.cmake @@ -1,5 +1,5 @@ ############################################################################## -# Copyright (c) 2016-19, Lawrence Livermore National Security, LLC and CHAI +# Copyright (c) 2016-20, Lawrence Livermore National Security, LLC and CHAI # project contributors. See the COPYRIGHT file for details. # # SPDX-License-Identifier: BSD-3-Clause diff --git a/docs/CMakeLists.txt b/docs/CMakeLists.txt index c50020b5..34a2ced2 100644 --- a/docs/CMakeLists.txt +++ b/docs/CMakeLists.txt @@ -1,5 +1,5 @@ ############################################################################## -# Copyright (c) 2016-19, Lawrence Livermore National Security, LLC and CHAI +# Copyright (c) 2016-20, Lawrence Livermore National Security, LLC and CHAI # project contributors. See the COPYRIGHT file for details. # # SPDX-License-Identifier: BSD-3-Clause diff --git a/docs/doxygen/CMakeLists.txt b/docs/doxygen/CMakeLists.txt index ac0fc239..c53655c2 100644 --- a/docs/doxygen/CMakeLists.txt +++ b/docs/doxygen/CMakeLists.txt @@ -1,5 +1,5 @@ ############################################################################## -# Copyright (c) 2016-19, Lawrence Livermore National Security, LLC and CHAI +# Copyright (c) 2016-20, Lawrence Livermore National Security, LLC and CHAI # project contributors. See the COPYRIGHT file for details. # # SPDX-License-Identifier: BSD-3-Clause diff --git a/docs/doxygen/Doxyfile.in b/docs/doxygen/Doxyfile.in index 1bb28a90..64c68ce8 100644 --- a/docs/doxygen/Doxyfile.in +++ b/docs/doxygen/Doxyfile.in @@ -1,5 +1,5 @@ ############################################################################## -# Copyright (c) 2016-19, Lawrence Livermore National Security, LLC and CHAI +# Copyright (c) 2016-20, Lawrence Livermore National Security, LLC and CHAI # project contributors. See the COPYRIGHT file for details. # # SPDX-License-Identifier: BSD-3-Clause diff --git a/docs/sphinx/CMakeLists.txt b/docs/sphinx/CMakeLists.txt index d9f4aed8..acc991ab 100644 --- a/docs/sphinx/CMakeLists.txt +++ b/docs/sphinx/CMakeLists.txt @@ -1,5 +1,5 @@ ############################################################################## -# Copyright (c) 2016-19, Lawrence Livermore National Security, LLC and CHAI +# Copyright (c) 2016-20, Lawrence Livermore National Security, LLC and CHAI # project contributors. See the COPYRIGHT file for details. # # SPDX-License-Identifier: BSD-3-Clause diff --git a/docs/sphinx/conf.py b/docs/sphinx/conf.py index 404b4881..42274274 100644 --- a/docs/sphinx/conf.py +++ b/docs/sphinx/conf.py @@ -1,5 +1,5 @@ ############################################################################## -# Copyright (c) 2016-19, Lawrence Livermore National Security, LLC and CHAI +# Copyright (c) 2016-20, Lawrence Livermore National Security, LLC and CHAI # project contributors. See the COPYRIGHT file for details. # # SPDX-License-Identifier: BSD-3-Clause diff --git a/docs/sphinx/conf.py.in b/docs/sphinx/conf.py.in index de241adb..a8df17f1 100644 --- a/docs/sphinx/conf.py.in +++ b/docs/sphinx/conf.py.in @@ -1,5 +1,5 @@ ############################################################################## -# Copyright (c) 2016-19, Lawrence Livermore National Security, LLC and CHAI +# Copyright (c) 2016-20, Lawrence Livermore National Security, LLC and CHAI # project contributors. See the COPYRIGHT file for details. # # SPDX-License-Identifier: BSD-3-Clause diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt index b4c83bf8..09397813 100644 --- a/examples/CMakeLists.txt +++ b/examples/CMakeLists.txt @@ -1,5 +1,5 @@ ############################################################################## -# Copyright (c) 2016-19, Lawrence Livermore National Security, LLC and CHAI +# Copyright (c) 2016-20, Lawrence Livermore National Security, LLC and CHAI # project contributors. See the COPYRIGHT file for details. # # SPDX-License-Identifier: BSD-3-Clause diff --git a/examples/chai-umpire-allocators.cpp b/examples/chai-umpire-allocators.cpp index 15fa4fa9..fd999f4f 100644 --- a/examples/chai-umpire-allocators.cpp +++ b/examples/chai-umpire-allocators.cpp @@ -1,5 +1,5 @@ ////////////////////////////////////////////////////////////////////////////// -// Copyright (c) 2016-19, Lawrence Livermore National Security, LLC and CHAI +// Copyright (c) 2016-20, Lawrence Livermore National Security, LLC and CHAI // project contributors. See the COPYRIGHT file for details. // // SPDX-License-Identifier: BSD-3-Clause diff --git a/examples/ex1.cpp b/examples/ex1.cpp index fdec4b3b..39a75e5d 100644 --- a/examples/ex1.cpp +++ b/examples/ex1.cpp @@ -1,5 +1,5 @@ ////////////////////////////////////////////////////////////////////////////// -// Copyright (c) 2016-19, Lawrence Livermore National Security, LLC and CHAI +// Copyright (c) 2016-20, Lawrence Livermore National Security, LLC and CHAI // project contributors. See the COPYRIGHT file for details. // // SPDX-License-Identifier: BSD-3-Clause diff --git a/examples/example.cpp b/examples/example.cpp index e3405160..8772ba20 100644 --- a/examples/example.cpp +++ b/examples/example.cpp @@ -1,5 +1,5 @@ ////////////////////////////////////////////////////////////////////////////// -// Copyright (c) 2016-19, Lawrence Livermore National Security, LLC and CHAI +// Copyright (c) 2016-20, Lawrence Livermore National Security, LLC and CHAI // project contributors. See the COPYRIGHT file for details. // // SPDX-License-Identifier: BSD-3-Clause diff --git a/examples/pinned.cpp b/examples/pinned.cpp index a1e13e62..c819b725 100644 --- a/examples/pinned.cpp +++ b/examples/pinned.cpp @@ -1,5 +1,5 @@ ////////////////////////////////////////////////////////////////////////////// -// Copyright (c) 2016-19, Lawrence Livermore National Security, LLC and CHAI +// Copyright (c) 2016-20, Lawrence Livermore National Security, LLC and CHAI // project contributors. See the COPYRIGHT file for details. // // SPDX-License-Identifier: BSD-3-Clause diff --git a/scripts/apply-license-info.sh b/scripts/apply-license-info.sh index 99f7aea3..cb56e08c 100755 --- a/scripts/apply-license-info.sh +++ b/scripts/apply-license-info.sh @@ -1,6 +1,6 @@ #!/usr/bin/env zsh ############################################################################## -# Copyright (c) 2016-19, Lawrence Livermore National Security, LLC and CHAI +# Copyright (c) 2016-20, Lawrence Livermore National Security, LLC and CHAI # project contributors. See the COPYRIGHT file for details. # # SPDX-License-Identifier: BSD-3-Clause diff --git a/scripts/format-source.sh b/scripts/format-source.sh index 0206cfbe..91081e1a 100755 --- a/scripts/format-source.sh +++ b/scripts/format-source.sh @@ -1,6 +1,6 @@ #!/usr/bin/env bash ############################################################################## -# Copyright (c) 2016-19, Lawrence Livermore National Security, LLC and CHAI +# Copyright (c) 2016-20, Lawrence Livermore National Security, LLC and CHAI # project contributors. See the COPYRIGHT file for details. # # SPDX-License-Identifier: BSD-3-Clause diff --git a/scripts/license.txt b/scripts/license.txt index 5b1c9027..23a065a7 100644 --- a/scripts/license.txt +++ b/scripts/license.txt @@ -1,4 +1,4 @@ -Copyright (c) 2016-19, Lawrence Livermore National Security, LLC and CHAI +Copyright (c) 2016-20, Lawrence Livermore National Security, LLC and CHAI project contributors. See the COPYRIGHT file for details. SPDX-License-Identifier: BSD-3-Clause diff --git a/scripts/make_release_tarball.sh b/scripts/make_release_tarball.sh index 6258c5b8..8cf0f40f 100755 --- a/scripts/make_release_tarball.sh +++ b/scripts/make_release_tarball.sh @@ -1,6 +1,6 @@ #!/bin/bash ############################################################################## -# Copyright (c) 2016-19, Lawrence Livermore National Security, LLC and CHAI +# Copyright (c) 2016-20, Lawrence Livermore National Security, LLC and CHAI # project contributors. See the COPYRIGHT file for details. # # SPDX-License-Identifier: BSD-3-Clause diff --git a/scripts/travis/build_and_test.sh b/scripts/travis/build_and_test.sh index 44eb570c..c96efaa5 100755 --- a/scripts/travis/build_and_test.sh +++ b/scripts/travis/build_and_test.sh @@ -1,6 +1,6 @@ #!/bin/bash ############################################################################## -# Copyright (c) 2016-19, Lawrence Livermore National Security, LLC and CHAI +# Copyright (c) 2016-20, Lawrence Livermore National Security, LLC and CHAI # project contributors. See the COPYRIGHT file for details. # # SPDX-License-Identifier: BSD-3-Clause diff --git a/scripts/travis/install_llvm.sh b/scripts/travis/install_llvm.sh index 2b6c69fd..499f3a57 100755 --- a/scripts/travis/install_llvm.sh +++ b/scripts/travis/install_llvm.sh @@ -1,5 +1,5 @@ ############################################################################## -# Copyright (c) 2016-19, Lawrence Livermore National Security, LLC and CHAI +# Copyright (c) 2016-20, Lawrence Livermore National Security, LLC and CHAI # project contributors. See the COPYRIGHT file for details. # # SPDX-License-Identifier: BSD-3-Clause diff --git a/scripts/update-copyright-year.sh b/scripts/update-copyright-year.sh index 6b411e07..a42a1bf6 100755 --- a/scripts/update-copyright-year.sh +++ b/scripts/update-copyright-year.sh @@ -14,7 +14,7 @@ RED="\033[1;31m" GREEN="\033[1;32m" NOCOLOR="\033[0m" -files_no_license=$(grep -l '2016,' \ +files_no_license=$(grep -l '2016-19,' \ benchmarks/**/*(^/) \ cmake/**/*(^/) \ docs/**/*~*rst(^/)\ @@ -28,7 +28,7 @@ if [ $files_no_license ]; then print "${RED} [!] Some files need copyright year updating: ${NOCOLOR}" echo "${files_no_license}" - echo ${files_no_license} | xargs sed -i '' 's/2016,/2016-2018,/' + echo ${files_no_license} | xargs sed -i '' 's/2016-19,/2016-20,/' print "${GREEN} [Ok] Copyright years updated." diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 2b18fdec..8536f824 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -1,5 +1,5 @@ ############################################################################## -# Copyright (c) 2016-19, Lawrence Livermore National Security, LLC and CHAI +# Copyright (c) 2016-20, Lawrence Livermore National Security, LLC and CHAI # project contributors. See the COPYRIGHT file for details. # # SPDX-License-Identifier: BSD-3-Clause diff --git a/src/chai/ArrayManager.cpp b/src/chai/ArrayManager.cpp index 80e5be84..11e6b394 100644 --- a/src/chai/ArrayManager.cpp +++ b/src/chai/ArrayManager.cpp @@ -1,5 +1,5 @@ ////////////////////////////////////////////////////////////////////////////// -// Copyright (c) 2016-19, Lawrence Livermore National Security, LLC and CHAI +// Copyright (c) 2016-20, Lawrence Livermore National Security, LLC and CHAI // project contributors. See the COPYRIGHT file for details. // // SPDX-License-Identifier: BSD-3-Clause diff --git a/src/chai/ArrayManager.hpp b/src/chai/ArrayManager.hpp index 9353fe1f..5328bae3 100644 --- a/src/chai/ArrayManager.hpp +++ b/src/chai/ArrayManager.hpp @@ -1,5 +1,5 @@ ////////////////////////////////////////////////////////////////////////////// -// Copyright (c) 2016-19, Lawrence Livermore National Security, LLC and CHAI +// Copyright (c) 2016-20, Lawrence Livermore National Security, LLC and CHAI // project contributors. See the COPYRIGHT file for details. // // SPDX-License-Identifier: BSD-3-Clause diff --git a/src/chai/ArrayManager.inl b/src/chai/ArrayManager.inl index 7ac78a63..e1af96b2 100644 --- a/src/chai/ArrayManager.inl +++ b/src/chai/ArrayManager.inl @@ -1,5 +1,5 @@ ////////////////////////////////////////////////////////////////////////////// -// Copyright (c) 2016-19, Lawrence Livermore National Security, LLC and CHAI +// Copyright (c) 2016-20, Lawrence Livermore National Security, LLC and CHAI // project contributors. See the COPYRIGHT file for details. // // SPDX-License-Identifier: BSD-3-Clause diff --git a/src/chai/CMakeLists.txt b/src/chai/CMakeLists.txt index e77f09ab..7491f479 100644 --- a/src/chai/CMakeLists.txt +++ b/src/chai/CMakeLists.txt @@ -1,5 +1,5 @@ ############################################################################## -# Copyright (c) 2016-19, Lawrence Livermore National Security, LLC and CHAI +# Copyright (c) 2016-20, Lawrence Livermore National Security, LLC and CHAI # project contributors. See the COPYRIGHT file for details. # # SPDX-License-Identifier: BSD-3-Clause diff --git a/src/chai/ChaiMacros.hpp b/src/chai/ChaiMacros.hpp index 6fd48d7d..a9125715 100644 --- a/src/chai/ChaiMacros.hpp +++ b/src/chai/ChaiMacros.hpp @@ -1,5 +1,5 @@ ////////////////////////////////////////////////////////////////////////////// -// Copyright (c) 2016-19, Lawrence Livermore National Security, LLC and CHAI +// Copyright (c) 2016-20, Lawrence Livermore National Security, LLC and CHAI // project contributors. See the COPYRIGHT file for details. // // SPDX-License-Identifier: BSD-3-Clause diff --git a/src/chai/ExecutionSpaces.hpp b/src/chai/ExecutionSpaces.hpp index dfc7b773..43346160 100644 --- a/src/chai/ExecutionSpaces.hpp +++ b/src/chai/ExecutionSpaces.hpp @@ -1,5 +1,5 @@ ////////////////////////////////////////////////////////////////////////////// -// Copyright (c) 2016-19, Lawrence Livermore National Security, LLC and CHAI +// Copyright (c) 2016-20, Lawrence Livermore National Security, LLC and CHAI // project contributors. See the COPYRIGHT file for details. // // SPDX-License-Identifier: BSD-3-Clause diff --git a/src/chai/ManagedArray.hpp b/src/chai/ManagedArray.hpp index 5af7cec1..4a6aa3e1 100644 --- a/src/chai/ManagedArray.hpp +++ b/src/chai/ManagedArray.hpp @@ -1,5 +1,5 @@ ////////////////////////////////////////////////////////////////////////////// -// Copyright (c) 2016-19, Lawrence Livermore National Security, LLC and CHAI +// Copyright (c) 2016-20, Lawrence Livermore National Security, LLC and CHAI // project contributors. See the COPYRIGHT file for details. // // SPDX-License-Identifier: BSD-3-Clause diff --git a/src/chai/ManagedArray.inl b/src/chai/ManagedArray.inl index 09e84a52..60413c04 100644 --- a/src/chai/ManagedArray.inl +++ b/src/chai/ManagedArray.inl @@ -1,5 +1,5 @@ ////////////////////////////////////////////////////////////////////////////// -// Copyright (c) 2016-19, Lawrence Livermore National Security, LLC and CHAI +// Copyright (c) 2016-20, Lawrence Livermore National Security, LLC and CHAI // project contributors. See the COPYRIGHT file for details. // // SPDX-License-Identifier: BSD-3-Clause diff --git a/src/chai/ManagedArray_thin.inl b/src/chai/ManagedArray_thin.inl index 2a64b187..8a63e2b6 100644 --- a/src/chai/ManagedArray_thin.inl +++ b/src/chai/ManagedArray_thin.inl @@ -1,5 +1,5 @@ ////////////////////////////////////////////////////////////////////////////// -// Copyright (c) 2016-19, Lawrence Livermore National Security, LLC and CHAI +// Copyright (c) 2016-20, Lawrence Livermore National Security, LLC and CHAI // project contributors. See the COPYRIGHT file for details. // // SPDX-License-Identifier: BSD-3-Clause diff --git a/src/chai/PointerRecord.hpp b/src/chai/PointerRecord.hpp index 49c0ce5e..9ff9c372 100644 --- a/src/chai/PointerRecord.hpp +++ b/src/chai/PointerRecord.hpp @@ -1,5 +1,5 @@ ////////////////////////////////////////////////////////////////////////////// -// Copyright (c) 2016-19, Lawrence Livermore National Security, LLC and CHAI +// Copyright (c) 2016-20, Lawrence Livermore National Security, LLC and CHAI // project contributors. See the COPYRIGHT file for details. // // SPDX-License-Identifier: BSD-3-Clause diff --git a/src/chai/Types.hpp b/src/chai/Types.hpp index aeb16bf7..0164a2e0 100644 --- a/src/chai/Types.hpp +++ b/src/chai/Types.hpp @@ -1,5 +1,5 @@ ////////////////////////////////////////////////////////////////////////////// -// Copyright (c) 2016-19, Lawrence Livermore National Security, LLC and CHAI +// Copyright (c) 2016-20, Lawrence Livermore National Security, LLC and CHAI // project contributors. See the COPYRIGHT file for details. // // SPDX-License-Identifier: BSD-3-Clause diff --git a/src/chai/chai-config.cmake.in b/src/chai/chai-config.cmake.in index 1cc6198f..4fc7af6a 100644 --- a/src/chai/chai-config.cmake.in +++ b/src/chai/chai-config.cmake.in @@ -1,5 +1,5 @@ ############################################################################## -# Copyright (c) 2016-19, Lawrence Livermore National Security, LLC and CHAI +# Copyright (c) 2016-20, Lawrence Livermore National Security, LLC and CHAI # project contributors. See the COPYRIGHT file for details. # # SPDX-License-Identifier: BSD-3-Clause diff --git a/src/chai/config.hpp.in b/src/chai/config.hpp.in index 4da4a8e3..26961f34 100644 --- a/src/chai/config.hpp.in +++ b/src/chai/config.hpp.in @@ -1,5 +1,5 @@ ////////////////////////////////////////////////////////////////////////////// -// Copyright (c) 2016-19, Lawrence Livermore National Security, LLC and CHAI +// Copyright (c) 2016-20, Lawrence Livermore National Security, LLC and CHAI // project contributors. See the COPYRIGHT file for details. // // SPDX-License-Identifier: BSD-3-Clause diff --git a/src/util/forall.hpp b/src/util/forall.hpp index ad0ce770..38a21a5a 100644 --- a/src/util/forall.hpp +++ b/src/util/forall.hpp @@ -1,5 +1,5 @@ ////////////////////////////////////////////////////////////////////////////// -// Copyright (c) 2016-19, Lawrence Livermore National Security, LLC and CHAI +// Copyright (c) 2016-20, Lawrence Livermore National Security, LLC and CHAI // project contributors. See the COPYRIGHT file for details. // // SPDX-License-Identifier: BSD-3-Clause diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index cc1c8705..b2d1cba5 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -1,5 +1,5 @@ ############################################################################## -# Copyright (c) 2016-19, Lawrence Livermore National Security, LLC and CHAI +# Copyright (c) 2016-20, Lawrence Livermore National Security, LLC and CHAI # project contributors. See the COPYRIGHT file for details. # # SPDX-License-Identifier: BSD-3-Clause diff --git a/tests/integration/CMakeLists.txt b/tests/integration/CMakeLists.txt index 21e2a1e3..6825dfeb 100644 --- a/tests/integration/CMakeLists.txt +++ b/tests/integration/CMakeLists.txt @@ -1,5 +1,5 @@ ############################################################################## -# Copyright (c) 2016-19, Lawrence Livermore National Security, LLC and CHAI +# Copyright (c) 2016-20, Lawrence Livermore National Security, LLC and CHAI # project contributors. See the COPYRIGHT file for details. # # SPDX-License-Identifier: BSD-3-Clause diff --git a/tests/integration/managed_array_tests.cpp b/tests/integration/managed_array_tests.cpp index 078b0d94..957c8fe9 100644 --- a/tests/integration/managed_array_tests.cpp +++ b/tests/integration/managed_array_tests.cpp @@ -1,5 +1,5 @@ ////////////////////////////////////////////////////////////////////////////// -// Copyright (c) 2016-19, Lawrence Livermore National Security, LLC and CHAI +// Copyright (c) 2016-20, Lawrence Livermore National Security, LLC and CHAI // project contributors. See the COPYRIGHT file for details. // // SPDX-License-Identifier: BSD-3-Clause diff --git a/tests/integration/raja-chai-nested.cpp b/tests/integration/raja-chai-nested.cpp index 5caee295..664f470a 100644 --- a/tests/integration/raja-chai-nested.cpp +++ b/tests/integration/raja-chai-nested.cpp @@ -1,5 +1,5 @@ //~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~// -// Copyright (c) 2016-19, Lawrence Livermore National Security, LLC +// Copyright (c) 2016-20, Lawrence Livermore National Security, LLC // and RAJA project contributors. See the RAJA/COPYRIGHT file for details. // // SPDX-License-Identifier: (BSD-3-Clause) diff --git a/tests/integration/raja-chai-tests.cpp b/tests/integration/raja-chai-tests.cpp index c1e9416c..a5cd7c87 100644 --- a/tests/integration/raja-chai-tests.cpp +++ b/tests/integration/raja-chai-tests.cpp @@ -1,5 +1,5 @@ //~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~// -// Copyright (c) 2016-19, Lawrence Livermore National Security, LLC +// Copyright (c) 2016-20, Lawrence Livermore National Security, LLC // and RAJA project contributors. See the RAJA/COPYRIGHT file for details. // // SPDX-License-Identifier: (BSD-3-Clause) diff --git a/tests/unit/CMakeLists.txt b/tests/unit/CMakeLists.txt index d7939e65..2438289e 100644 --- a/tests/unit/CMakeLists.txt +++ b/tests/unit/CMakeLists.txt @@ -1,5 +1,5 @@ ############################################################################## -# Copyright (c) 2016-19, Lawrence Livermore National Security, LLC and CHAI +# Copyright (c) 2016-20, Lawrence Livermore National Security, LLC and CHAI # project contributors. See the COPYRIGHT file for details. # # SPDX-License-Identifier: BSD-3-Clause diff --git a/tests/unit/array_manager_unit_tests.cpp b/tests/unit/array_manager_unit_tests.cpp index 1c212f21..d03b0577 100644 --- a/tests/unit/array_manager_unit_tests.cpp +++ b/tests/unit/array_manager_unit_tests.cpp @@ -1,5 +1,5 @@ ////////////////////////////////////////////////////////////////////////////// -// Copyright (c) 2016-19, Lawrence Livermore National Security, LLC and CHAI +// Copyright (c) 2016-20, Lawrence Livermore National Security, LLC and CHAI // project contributors. See the COPYRIGHT file for details. // // SPDX-License-Identifier: BSD-3-Clause diff --git a/tests/unit/managed_array_unit_tests.cpp b/tests/unit/managed_array_unit_tests.cpp index 50fb1696..0d3e3dae 100644 --- a/tests/unit/managed_array_unit_tests.cpp +++ b/tests/unit/managed_array_unit_tests.cpp @@ -1,5 +1,5 @@ ////////////////////////////////////////////////////////////////////////////// -// Copyright (c) 2016-19, Lawrence Livermore National Security, LLC and CHAI +// Copyright (c) 2016-20, Lawrence Livermore National Security, LLC and CHAI // project contributors. See the COPYRIGHT file for details. // // SPDX-License-Identifier: BSD-3-Clause