Merge pull request #324 from LLNL/v0.15.1-rc

v0.15.1 RC

Author: adayton1

CMakeLists.txt

@@ -17,7 +17,7 @@ cmake_minimum_required(VERSION 3.23)
 
 project(CARE
         LANGUAGES C CXX
-        VERSION 0.15.0)
+        VERSION 0.15.1)
 
 include(${PROJECT_SOURCE_DIR}/cmake/Setup.cmake)
 

README.md

@@ -5,7 +5,7 @@
 [comment]: # (SPDX-License-Identifier: BSD-3-Clause)
 [comment]: # (#################################################################)
 
-# CARE v0.15.0
+# CARE v0.15.1
 
 CARE: CHAI and RAJA Extensions
 ===============================
@@ -21,7 +21,7 @@ cmake ../   # May need to pass -DCMAKE_INSTALL_PREFIX=/path/to/install/in if the
 make -j install
 ```
 
-If desired, external libraries can be used instead of submodules. For example, an external CHAI can be specified with `-DCHAI_DIR=<path to CHAI install directory or directory containing chai-config.cmake>`. Note that if using an external CHAI, it must be configured with `-DENABLE_PICK=ON -DENABLE_PINNED=ON`.
+If desired, external libraries can be used instead of submodules. For example, an external CHAI can be specified with `-DCHAI_DIR=<path to CHAI install directory or directory containing chai-config.cmake>`. Note that if using an external CHAI, it must be configured with `-DENABLE_PINNED=ON`.
 
 To build with CUDA support, use `-DENABLE_CUDA=ON -DCUDA_TOOLKIT_ROOT_DIR=/path/to/cuda/toolkit`. If using external libraries, note that Umpire, RAJA, and CHAI must also be configured with those options.
 

RELEASE_NOTES.md

@@ -12,6 +12,28 @@ in this file.
 
 The format of this file is based on [Keep a Changelog](http://keepachangelog.com/en/1.0.0/).
 
+## [Version 0.15.1] - Release date 2025-04-07
+
+### Added
+- Added support for RAJA MultiReducers (Min/Max/Sum).
+
+### Changed
+- Changed RAJA reduce policy for CUDA to RAJA::cuda\_reduce\_atomic.
+- Rearranged template parameters of care::sortKeyValueArrays (used by care::KeyValueSorter) for ease of use
+
+### Fixed
+- Only enable calls to cub::DeviceMergeSort when it is available (used by care::sortArray and care::KeyValueSorter when the type is not arithmetic)
+- Fixes inputs to [hip]cub::DeviceMergeSort::StableSortKeys (used by care::sortArray when the type is not arithmetic)
+- Avoids hardcoding one overload of care::sortArray to use [hip]cub::DeviceRadixSort
+- Fixes a case where care::sort\_uniq should not modify the input array
+- Miscellaneous fixes for care::host\_device\_map
+- Clarified documentation for care::BinarySearch
+- Added missing attributes to functions for building as a shared library on Windows
+- Moved helper function to be accessible when the loop fuser is disabled
+
+### Removed
+- Removed dead ENABLE\_PICK option (corresponding option has been removed from CHAI)
+
 ## [Version 0.15.0] - Release date 2025-03-20
 
 ### Added

cmake/SetupDependencies.cmake

@@ -83,7 +83,6 @@ if(NOT TARGET chai)
       if(NOT EXISTS ${PROJECT_SOURCE_DIR}/tpl/chai/CMakeLists.txt)
          message(FATAL_ERROR "CARE: CHAI submodule not initialized. Run 'git submodule update --init' in the git repository or set CHAI_DIR to use an external build of CHAI.")
       else()
-         set(CHAI_ENABLE_PICK ${ENABLE_PICK} CACHE BOOL "Enable picks/sets in chai::ManagedArray")
          set(CHAI_ENABLE_PINNED ${ENABLE_PINNED} CACHE BOOL "Enable pinned memory support in CHAI")
 
          set(CHAI_ENABLE_TESTS ${CARE_ENABLE_SUBMODULE_TESTS} CACHE BOOL "Enable CHAI tests")

cmake/SetupOptions.cmake

@@ -10,7 +10,6 @@ include(CMakeDependentOption)
 # Advanced configuration options
 # TODO: If these are disabled, the code will not compile or run correctly.
 #       Fix those issues so that these options are actually configurable.
-option(ENABLE_PICK "Enable pick and set methods on ManagedArrays" ON)
 option(ENABLE_PINNED "Enable pinned memory space" ON)
 option(CARE_ENABLE_PINNED_MEMORY_FOR_SCANS "Use pinned memory for scan lengths" ON)
 option(CARE_GPU_MEMORY_IS_ACCESSIBLE_ON_CPU "Allows default memory spaces for ZERO_COPY and PAGEABLE to be the GPU memory space" OFF)

docs/sphinx/conf.py

@@ -57,7 +57,7 @@
 # The short X.Y version.
 version = '0.15'
 # The full version, including alpha/beta/rc tags.
-release = '0.15.0'
+release = '0.15.1'
 
 # The language for content autogenerated by Sphinx. Refer to documentation
 # for a list of supported languages.

scripts/make_release_tarball.sh

@@ -8,7 +8,7 @@
 ##############################################################################
 
 TAR_CMD=gtar
-VERSION=0.15.0
+VERSION=0.15.1
 
 git archive --prefix=care-${VERSION}/ -o care-${VERSION}.tar HEAD 2> /dev/null
 

src/care/DebugPlugin.cpp

@@ -27,10 +27,6 @@ namespace care{
       // Prepare to record CHAI data
       if (CHAICallback::isActive()) {
          PluginData::clearActivePointers();
-
-#if defined(CARE_GPUCC) && defined(CARE_DEBUG)
-         GPUWatchpoint::setOrCheckWatchpoint<int>();
-#endif // defined(CARE_GPUCC) && defined(CARE_DEBUG)
       }
 #endif // !defined(CHAI_DISABLE_RM)
    }
@@ -57,10 +53,6 @@ namespace care{
 
       if (CHAICallback::isActive()) {			
          writeLoopData(space, PluginData::getFileName(), PluginData::getLineNumber());
-
-#if defined(CARE_GPUCC) && defined(CARE_DEBUG)
-      GPUWatchpoint::setOrCheckWatchpoint<int>();
-#endif // defined(CARE_GPUCC) && defined(CARE_DEBUG)
       }
 
       if (PluginData::isParallelContext()) {

src/care/DebugPlugin.h

@@ -8,6 +8,7 @@
 #ifndef _CARE_DebugPlugin_H_
 #define _CARE_DebugPlugin_H_
 
+#include "care/config.h"
 #include "RAJA/util/PluginStrategy.hpp"
 #include "chai/ExecutionSpaces.hpp"
 
@@ -18,7 +19,7 @@ namespace care{
       public:
          DebugPlugin() = default;
 
-         static void registerPlugin();
+         CARE_DLL_API static void registerPlugin();
 
          void preLaunch(const RAJA::util::PluginContext& p) override;
 

src/care/ExecutionSpace.h

@@ -31,19 +31,18 @@ namespace care {
    // CHAI and CARE are configured.
    extern CARE_DLL_API care::ExecutionSpace DEFAULT;
 } // namespace care
+
 namespace chai {
-   
    // the ZERO_COPY memory space. Typically PINNED memory, but may be a different space depending on
    // how CHAI and CARE are configured.
-   extern chai::ExecutionSpace ZERO_COPY;
+   extern CARE_DLL_API chai::ExecutionSpace ZERO_COPY;
    // the  PAGEABLE memory space. Typically UM, but may be a different space depending on how
    // CHAI and CARE are configured.
-   extern chai::ExecutionSpace PAGEABLE;
+   extern CARE_DLL_API chai::ExecutionSpace PAGEABLE;
    // the  DEFAULT memory space. Typically GPU for GPU platforms and CPU for CPU platforms, but may be a different space depending on how
    // CHAI and CARE are configured.
-   extern chai::ExecutionSpace DEFAULT;
-
-}
+   extern CARE_DLL_API chai::ExecutionSpace DEFAULT;
+} // namespace chai
 
 #endif // !defined(_CARE_EXECUTION_SPACE_H_)
 

src/care/GPUMacros.h

@@ -9,6 +9,7 @@
 #define _CARE_GPUMACROS_H_
 
 #include "chai/config.hpp"
+#include "care/config.h"
 
 #if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__)
 #define CARE_DEVICE_COMPILE
@@ -58,10 +59,10 @@
 #define gpuMemGetInfo cudaMemGetInfo
 
 #define gpuDeviceSetLimit cudaDeviceSetLimit
-#define gpuDeviceSynchronize cudaDeviceSynchronize
 #define gpuLimitStackSize cudaLimitStackSize
 #define gpuLimitMallocHeapSize cudaLimitMallocHeapSize
 
+#define gpuDeviceSynchronize cudaDeviceSynchronize
 #define gpuPeekAtLastError cudaPeekAtLastError
 
 #define gpuSuccess cudaSuccess
@@ -81,15 +82,21 @@
 #define gpuMemGetInfo hipMemGetInfo
 
 #define gpuDeviceSetLimit hipDeviceSetLimit
-#define gpuDeviceSynchronize hipDeviceSynchronize
 #define gpuLimitStackSize hipLimitStackSize
 #define gpuLimitMallocHeapSize hipLimitMallocHeapSize
 
+#define gpuDeviceSynchronize hipDeviceSynchronize
 #define gpuPeekAtLastError hipPeekAtLastError
 
 #define gpuSuccess hipSuccess
 
-#endif // end __HIPCC__ case
+#elif CARE_ENABLE_GPU_SIMULATION_MODE
+
+#define gpuMemcpyKind int
+#define gpuSimNoop()
+#define gpuPeekAtLastError gpuSimNoop
+
+#endif // #if defined(__CUDACC__) #elif defined(__HIPCC__)
 
 #endif // !defined(_CARE_GPUMACROS_H_)
 

src/care/KeyValueSorter_decl.h

@@ -55,11 +55,21 @@ using LocalKeyValueSorter = KeyValueSorter<KeyType, ValueType, Exec> ;
 ///                             have bugs!
 /// @return void
 ///////////////////////////////////////////////////////////////////////////
-template <typename KeyT, typename ValueT, typename Exec=RAJADeviceExec>
-void sortKeyValueArrays(host_device_ptr<KeyT> & keys,
-                        host_device_ptr<ValueT> & values,
-                        const size_t start, const size_t len,
-                        const bool noCopy=false) ;
+template <typename Exec, typename KeyT, typename ValueT>
+std::enable_if_t<std::is_arithmetic<typename CHAIDataGetter<KeyT, RAJADeviceExec>::raw_type>::value, void>
+sortKeyValueArrays(host_device_ptr<KeyT> & keys,
+                   host_device_ptr<ValueT> & values,
+                   const size_t start, const size_t len,
+                   const bool noCopy=false);
+
+#if defined(__HIPCC__) || (defined(__CUDACC__) && defined(CUB_MAJOR_VERSION) && defined(CUB_MINOR_VERSION) && (CUB_MAJOR_VERSION >= 2 || (CUB_MAJOR_VERSION == 1 && CUB_MINOR_VERSION >= 14)))
+template <typename Exec, typename KeyT, typename ValueT>
+std::enable_if_t<!std::is_arithmetic<typename CHAIDataGetter<KeyT, RAJADeviceExec>::raw_type>::value, void>
+sortKeyValueArrays(host_device_ptr<KeyT> & keys,
+                   host_device_ptr<ValueT> & values,
+                   const size_t start, const size_t len,
+                   const bool noCopy=false);
+#endif
 
 ///////////////////////////////////////////////////////////////////////////
 /// @author Benjamin Liu after Alan Dayton
@@ -358,7 +368,7 @@ class CARE_DLL_API KeyValueSorter<KeyType, ValueType, RAJADeviceExec> {
       /// TODO: add bounds checking
       ///////////////////////////////////////////////////////////////////////////
       void sort(const size_t start, const size_t len) {
-         sortKeyValueArrays<ValueType, KeyType, RAJADeviceExec>(m_values, m_keys, start, len, false);
+         sortKeyValueArrays<RAJADeviceExec>(m_values, m_keys, start, len, false);
       }
 
       ///////////////////////////////////////////////////////////////////////////
@@ -377,7 +387,7 @@ class CARE_DLL_API KeyValueSorter<KeyType, ValueType, RAJADeviceExec> {
       /// @return void
       ///////////////////////////////////////////////////////////////////////////
       void sort() {
-         sortKeyValueArrays<ValueType, KeyType, RAJADeviceExec>(m_values, m_keys, 0, m_len, true);
+         sortKeyValueArrays<RAJADeviceExec>(m_values, m_keys, 0, m_len, true);
       }
 
       ///////////////////////////////////////////////////////////////////////////
@@ -389,7 +399,7 @@ class CARE_DLL_API KeyValueSorter<KeyType, ValueType, RAJADeviceExec> {
       /// TODO: add bounds checking
       ///////////////////////////////////////////////////////////////////////////
       void sortByKey(const size_t start, const size_t len) {
-         sortKeyValueArrays<KeyType, ValueType, RAJADeviceExec>(m_keys, m_values, start, len, false);
+         sortKeyValueArrays<RAJADeviceExec>(m_keys, m_values, start, len, false);
       }
 
       ///////////////////////////////////////////////////////////////////////////
@@ -408,7 +418,7 @@ class CARE_DLL_API KeyValueSorter<KeyType, ValueType, RAJADeviceExec> {
       /// @return void
       ///////////////////////////////////////////////////////////////////////////
       void sortByKey() {
-         sortKeyValueArrays(m_keys, m_values, 0, m_len, true);
+         sortKeyValueArrays<RAJADeviceExec>(m_keys, m_values, 0, m_len, true);
       }
 
       ///////////////////////////////////////////////////////////////////////////
@@ -421,7 +431,7 @@ class CARE_DLL_API KeyValueSorter<KeyType, ValueType, RAJADeviceExec> {
       /// TODO: add bounds checking
       ///////////////////////////////////////////////////////////////////////////
       void stableSort(const size_t start, const size_t len) {
-         sortKeyValueArrays(m_values, m_keys, start, len, false);
+         sortKeyValueArrays<RAJADeviceExec>(m_values, m_keys, start, len, false);
       }
 
       ///////////////////////////////////////////////////////////////////////////
@@ -441,7 +451,7 @@ class CARE_DLL_API KeyValueSorter<KeyType, ValueType, RAJADeviceExec> {
       /// TODO: investigate whether radix device sort is a stable sort
       ///////////////////////////////////////////////////////////////////////////
       void stableSort() {
-         sortKeyValueArrays(m_values, m_keys, 0, m_len, true);
+         sortKeyValueArrays<RAJADeviceExec>(m_values, m_keys, 0, m_len, true);
       }
 
       ///////////////////////////////////////////////////////////////////////////
@@ -1181,11 +1191,12 @@ class CARE_DLL_API KeyValueSorter<KeyType, ValueType, RAJA::seq_exec> {
 #endif // !CARE_ENABLE_GPU_SIMULATION_MODE
 
 
+// Return the keys for each KVS where their values are the same
 #ifdef CARE_PARALLEL_DEVICE
 template <typename KeyType, typename ValueType>
 void IntersectKeyValueSorters(RAJADeviceExec exec, KeyValueSorter<KeyType, ValueType, RAJADeviceExec> sorter1, int size1,
                               KeyValueSorter<KeyType, ValueType, RAJADeviceExec> sorter2, int size2,
-                              host_device_ptr<int> &matches1, host_device_ptr<int>& matches2,
+                              host_device_ptr<KeyType> &matches1, host_device_ptr<KeyType>& matches2,
                               int & numMatches) ;
 #endif // defined(CARE_PARALLEL_DEVICE)
 
@@ -1197,7 +1208,7 @@ template <typename KeyType, typename ValueType>
 void IntersectKeyValueSorters(RAJA::seq_exec exec, 
                               KeyValueSorter<KeyType, ValueType, RAJA::seq_exec> sorter1, int size1,
                               KeyValueSorter<KeyType, ValueType, RAJA::seq_exec> sorter2, int size2,
-                              host_device_ptr<int> &matches1, host_device_ptr<int>& matches2, int & numMatches) ;
+                              host_device_ptr<KeyType> &matches1, host_device_ptr<KeyType>& matches2, int & numMatches) ;
 
 } // namespace care