From d59cd4108f441e5c632c469280575d6ae77ff24e Mon Sep 17 00:00:00 2001 From: "Peter B. Robinson" Date: Thu, 20 Aug 2026 13:18:55 -0700 Subject: [PATCH] actual changes --- src/care/KeyValueSorter_decl.h | 275 ++++++++++++++++++++++++- src/care/KeyValueSorter_impl.h | 359 +++------------------------------ src/care/LoopFuser.cpp | 5 +- src/care/algorithm_decl.h | 304 ++++++++++++++++++++++++++++ src/care/algorithm_impl.h | 296 --------------------------- src/care/care_inst.h | 171 ---------------- 6 files changed, 606 insertions(+), 804 deletions(-) diff --git a/src/care/KeyValueSorter_decl.h b/src/care/KeyValueSorter_decl.h index 514ebf34..038c135e 100644 --- a/src/care/KeyValueSorter_decl.h +++ b/src/care/KeyValueSorter_decl.h @@ -16,6 +16,19 @@ #include "care/scan.h" +// Other library headers +#ifdef CARE_GPUCC +#if defined(__CUDACC__) +#include "cub/cub.cuh" +#undef CUB_NS_POSTFIX +#undef CUB_NS_PREFIX +#endif + +#if defined(__HIPCC__) +#include "hipcub/hipcub.hpp" +#endif +#endif + #include // For std::move namespace care { @@ -42,6 +55,39 @@ class CARE_KEY_VALUE_SORTER_DLL_API KeyValueSorter; template using LocalKeyValueSorter = KeyValueSorter ; +template +inline bool cmpKeys(KeyValueType const & left, KeyValueType const & right); + +template +inline bool cmpKeysThenValues(KeyValueType const & left, KeyValueType const & right); + +namespace detail { + template + CARE_INLINE void stableSortKeyValuePairs(host_device_ptr & keys, + host_device_ptr & values, + const size_t len, + const size_t start = 0) { + + host_device_ptr<_kv> keyValues(len); + + CARE_SEQUENTIAL_LOOP(i, 0, (int) len) { + keyValues[i].key = keys[i+start]; + keyValues[i].value = values[i+start]; + } CARE_SEQUENTIAL_LOOP_END + + CHAIDataGetter<_kv, RAJA::seq_exec> getter {}; + _kv * rawData = getter.getRawArrayData(keyValues); + std::stable_sort(rawData, rawData + len, cmpKeys<_kv>); + + CARE_SEQUENTIAL_LOOP(i, 0, (int) len) { + keys[i+start] = keyValues[i].key; + values[i+start] = keyValues[i].value; + } CARE_SEQUENTIAL_LOOP_END + + keyValues.free(); + } +} // namespace detail + /////////////////////////////////////////////////////////////////////////// @@ -60,19 +106,230 @@ using LocalKeyValueSorter = KeyValueSorter ; /////////////////////////////////////////////////////////////////////////// template std::enable_if_t::raw_type>::value, void> -sortKeyValueArrays(host_device_ptr & keys, - host_device_ptr & values, - const size_t start, const size_t len, - const bool noCopy=false); +inline sortKeyValueArrays(host_device_ptr & keys, + host_device_ptr & values, + const size_t start, const size_t len, + const bool noCopy=false) +{ + bool _noCopy ; + if (noCopy && start > 0) { + printf("[CARE] Warning: sortKeyValueArrays. noCopy should not be set if start > 0 (%d)\n", (int)start); + _noCopy = false; + } + else { + _noCopy = noCopy; + } + + if constexpr (std::is_same_v) { + detail::stableSortKeyValuePairs(keys, values, len, start); + } + else { + // TODO openMP parallel implementation +#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))) + + // Allocate space for the result + host_device_ptr keyResult{len}; + host_device_ptr valueResult{len}; + + // Get the raw data to pass to cub + CHAIDataGetter keyGetter {}; + CHAIDataGetter valueGetter {}; + + auto * rawKeyData = keyGetter.getRawArrayData(keys) + start; + auto * rawValueData = valueGetter.getRawArrayData(values) + start; + + auto * rawKeyResult = keyGetter.getRawArrayData(keyResult); + auto * rawValueResult = valueGetter.getRawArrayData(valueResult); + + // Get the temp storage length + char * d_temp_storage = nullptr; + size_t temp_storage_bytes = 0; + + // When called with a nullptr for temp storage, this returns how much + // temp storage should be allocated. + if (len > 0) { +#if defined(__CUDACC__) + cub::DeviceRadixSort::SortPairs((void *)d_temp_storage, temp_storage_bytes, + rawKeyData, rawKeyResult, + rawValueData, rawValueResult, + len); +#elif defined(__HIPCC__) + hipcub::DeviceRadixSort::SortPairs((void *)d_temp_storage, temp_storage_bytes, + rawKeyData, rawKeyResult, + rawValueData, rawValueResult, + len); +#endif + } + + // Allocate the temp storage and get raw data to pass to cub + host_device_ptr tmpManaged {temp_storage_bytes}; + + CHAIDataGetter charGetter {}; + d_temp_storage = charGetter.getRawArrayData(tmpManaged); + + // Now sort + if (len > 0) { +#if defined(CHAI_THIN_GPU_ALLOCATE) + chai::ArrayManager::getInstance()->setExecutionSpace(chai::GPU); +#endif +#if defined(__CUDACC__) + cub::DeviceRadixSort::SortPairs((void *)d_temp_storage, temp_storage_bytes, + rawKeyData, rawKeyResult, + rawValueData, rawValueResult, + len); +#elif defined(__HIPCC__) + hipcub::DeviceRadixSort::SortPairs((void *)d_temp_storage, temp_storage_bytes, + rawKeyData, rawKeyResult, + rawValueData, rawValueResult, + len); +#endif + +#if defined(CHAI_THIN_GPU_ALLOCATE) + chai::ArrayManager::getInstance()->setExecutionSpace(chai::NONE); +#endif + + tmpManaged.free(); + } + + // Get the result + if (_noCopy) { + if (len > 0) { + keys.free(); + values.free(); + } + + keys = keyResult; + values = valueResult; + } + else { + CARE_STREAM_LOOP(i, 0, len) { + keys[i+start] = keyResult[i]; + values[i+start] = valueResult[i]; + } CARE_STREAM_LOOP_END + + if (len > 0) { + keyResult.free(); + valueResult.free(); + } + } + +#else // defined(CARE_GPUCC) + detail::stableSortKeyValuePairs(keys, values, len, start); +#endif // defined(CARE_GPUCC) + } + +} template std::enable_if_t::raw_type>::value, void> -sortKeyValueArrays(host_device_ptr & keys, - host_device_ptr & values, - const size_t start, const size_t len, - const bool noCopy=false); +inline sortKeyValueArrays(host_device_ptr & keys, + host_device_ptr & values, + const size_t start, const size_t len, + const bool noCopy=false) +{ + bool _noCopy ; + if (noCopy && start > 0) { + printf("[CARE] Warning: sortKeyValueArrays. noCopy should not be set if start > 0 (%d)\n", (int)start); + _noCopy = false; + } + else { + _noCopy = noCopy; + } + + if constexpr (std::is_same_v) { + detail::stableSortKeyValuePairs(keys, values, len, start); + } + else { + // TODO openMP parallel implementation +#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))) + + // Allocate space for the result + host_device_ptr keyResult{len}; + host_device_ptr valueResult{len}; + + // Get the raw data to pass to cub + CHAIDataGetter keyGetter {}; + CHAIDataGetter valueGetter {}; + + auto * rawKeyData = keyGetter.getRawArrayData(keys) + start; + auto * rawValueData = valueGetter.getRawArrayData(values) + start; + + auto * rawKeyResult = keyGetter.getRawArrayData(keyResult); + auto * rawValueResult = valueGetter.getRawArrayData(valueResult); + + using RawKeyType = std::remove_reference_t; + + auto custom_comparator = [] CARE_HOST_DEVICE (const RawKeyType& lhs, + const RawKeyType& rhs) { + return lhs < rhs; + }; + + // Get the temp storage length + char * d_temp_storage = nullptr; + size_t temp_storage_bytes = 0; + + // When called with a nullptr for temp storage, this returns how much + // temp storage should be allocated. + if (len > 0) { +#if defined(__CUDACC__) + cub::DeviceMergeSort::StableSortPairs((void *)d_temp_storage, temp_storage_bytes, + rawKeyData, + rawValueData, + len, custom_comparator); +#elif defined(__HIPCC__) + hipcub::DeviceMergeSort::StableSortPairs((void *)d_temp_storage, temp_storage_bytes, + rawKeyData, + rawValueData, + len, custom_comparator); +#endif + } + // Allocate the temp storage and get raw data to pass to cub + host_device_ptr tmpManaged {temp_storage_bytes}; + + CHAIDataGetter charGetter {}; + d_temp_storage = charGetter.getRawArrayData(tmpManaged); + + // Now sort + if (len > 0) { +#if defined(CHAI_THIN_GPU_ALLOCATE) + chai::ArrayManager::getInstance()->setExecutionSpace(chai::GPU); +#endif + +#if defined(__CUDACC__) + cub::DeviceMergeSort::StableSortPairs((void *)d_temp_storage, temp_storage_bytes, + rawKeyData, + rawValueData, + len, + custom_comparator); +#elif defined(__HIPCC__) + hipcub::DeviceMergeSort::StableSortPairs((void *)d_temp_storage, temp_storage_bytes, + rawKeyData, + rawValueData, + len, + custom_comparator); +#endif + +#if defined(CHAI_THIN_GPU_ALLOCATE) + chai::ArrayManager::getInstance()->setExecutionSpace(chai::NONE); +#endif + + tmpManaged.free(); + } + + // merge sort did an inplace sort, so the answer is already in keys and Values + if (len > 0) { + keyResult.free(); + valueResult.free(); + } + +#else // defined(CARE_GPUCC) + detail::stableSortKeyValuePairs(keys, values, len, start); +#endif // defined(CARE_GPUCC) + } + +} #if defined(CARE_PARALLEL_DEVICE) || CARE_ENABLE_GPU_SIMULATION_MODE /////////////////////////////////////////////////////////////////////////// @@ -617,7 +874,7 @@ class CARE_KEY_VALUE_SORTER_DLL_API KeyValueSorter positions(m_len+1); - care::exclusive_scan(RAJADeviceExec{}, isUnique, positions, m_len + 1, 0, false); + exclusive_scan(RAJADeviceExec{}, isUnique, positions, m_len + 1, 0, false); // Get the total number of unique elements int newSize = positions.pick(m_len); diff --git a/src/care/KeyValueSorter_impl.h b/src/care/KeyValueSorter_impl.h index b9039bdf..3ac475cc 100644 --- a/src/care/KeyValueSorter_impl.h +++ b/src/care/KeyValueSorter_impl.h @@ -39,302 +39,14 @@ namespace care { // TODO openMP parallel implementation -namespace detail { -template -CARE_INLINE void stableSortKeyValuePairs(host_device_ptr & keys, - host_device_ptr & values, - const size_t len, - const size_t start=0) -{ - host_device_ptr<_kv> keyValues(len); - - CARE_SEQUENTIAL_LOOP(i, 0, (int) len) { - keyValues[i].key = keys[i+start]; - keyValues[i].value = values[i+start]; - } CARE_SEQUENTIAL_LOOP_END - - CHAIDataGetter<_kv, RAJA::seq_exec> getter {}; - _kv * rawData = getter.getRawArrayData(keyValues); - std::stable_sort(rawData, rawData + len, cmpKeys<_kv>); - - CARE_SEQUENTIAL_LOOP(i, 0, (int) len) { - keys[i+start] = keyValues[i].key; - values[i+start] = keyValues[i].value; - } CARE_SEQUENTIAL_LOOP_END - - keyValues.free(); -} - -} // namespace detail - - -/////////////////////////////////////////////////////////////////////////// -/// @author Peter Robinson, Alan Dayton -/// @brief ManagedArray API for sorting paired key and value arrays -/// @param[in, out] keys - The array to sort -/// @param[in, out] values - The array that is sorted simultaneously -/// @param[in] start - The index to start sorting at -/// @param[in] len - The number of elements to sort -/// @param[in] noCopy - Whether or not to copy the result into the -/// original arrays or simply replace the -/// original arrays. Should be false if only -/// sorting part of the arrays or you will -/// have bugs! -/// @return void -/////////////////////////////////////////////////////////////////////////// -template -CARE_INLINE -std::enable_if_t::raw_type>::value, void> -sortKeyValueArrays(host_device_ptr & keys, - host_device_ptr & values, - const size_t start, const size_t len, - const bool noCopy) -{ - bool _noCopy ; - if (noCopy && start > 0) { - printf("[CARE] Warning: sortKeyValueArrays. noCopy should not be set if start > 0 (%d)\n", (int)start); - _noCopy = false; - } - else { - _noCopy = noCopy; - } - - if constexpr (std::is_same_v) { - detail::stableSortKeyValuePairs(keys, values, len, start); - } - else { - // TODO openMP parallel implementation -#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))) - - // Allocate space for the result - host_device_ptr keyResult{len}; - host_device_ptr valueResult{len}; - - // Get the raw data to pass to cub - CHAIDataGetter keyGetter {}; - CHAIDataGetter valueGetter {}; - - auto * rawKeyData = keyGetter.getRawArrayData(keys) + start; - auto * rawValueData = valueGetter.getRawArrayData(values) + start; - - auto * rawKeyResult = keyGetter.getRawArrayData(keyResult); - auto * rawValueResult = valueGetter.getRawArrayData(valueResult); - - // Get the temp storage length - char * d_temp_storage = nullptr; - size_t temp_storage_bytes = 0; - - // When called with a nullptr for temp storage, this returns how much - // temp storage should be allocated. - if (len > 0) { -#if defined(__CUDACC__) - cub::DeviceRadixSort::SortPairs((void *)d_temp_storage, temp_storage_bytes, - rawKeyData, rawKeyResult, - rawValueData, rawValueResult, - len); -#elif defined(__HIPCC__) - hipcub::DeviceRadixSort::SortPairs((void *)d_temp_storage, temp_storage_bytes, - rawKeyData, rawKeyResult, - rawValueData, rawValueResult, - len); -#endif - } - - // Allocate the temp storage and get raw data to pass to cub - host_device_ptr tmpManaged {temp_storage_bytes}; - - CHAIDataGetter charGetter {}; - d_temp_storage = charGetter.getRawArrayData(tmpManaged); - - // Now sort - if (len > 0) { -#if defined(CHAI_THIN_GPU_ALLOCATE) - chai::ArrayManager::getInstance()->setExecutionSpace(chai::GPU); -#endif - -#if defined(__CUDACC__) - cub::DeviceRadixSort::SortPairs((void *)d_temp_storage, temp_storage_bytes, - rawKeyData, rawKeyResult, - rawValueData, rawValueResult, - len); -#elif defined(__HIPCC__) - hipcub::DeviceRadixSort::SortPairs((void *)d_temp_storage, temp_storage_bytes, - rawKeyData, rawKeyResult, - rawValueData, rawValueResult, - len); -#endif - -#if defined(CHAI_THIN_GPU_ALLOCATE) - chai::ArrayManager::getInstance()->setExecutionSpace(chai::NONE); -#endif - - tmpManaged.free(); - } - - // Get the result - if (_noCopy) { - if (len > 0) { - keys.free(); - values.free(); - } - - keys = keyResult; - values = valueResult; - } - else { - CARE_STREAM_LOOP(i, 0, len) { - keys[i+start] = keyResult[i]; - values[i+start] = valueResult[i]; - } CARE_STREAM_LOOP_END - - if (len > 0) { - keyResult.free(); - valueResult.free(); - } - } - -#else // defined(CARE_GPUCC) - // fall back to an implementation that uses std::stable_sort on the host - detail::stableSortKeyValuePairs(keys, values, len, start); -#endif // defined(CARE_GPUCC) - } -} - -/////////////////////////////////////////////////////////////////////////// -/// @author Peter Robinson, Alan Dayton -/// @brief ManagedArray API for sorting paired key and value arrays -/// @param[in, out] keys - The array to sort -/// @param[in, out] values - The array that is sorted simultaneously -/// @param[in] start - The index to start sorting at -/// @param[in] len - The number of elements to sort -/// @param[in] noCopy - Whether or not to copy the result into the -/// original arrays or simply replace the -/// original arrays. Should be false if only -/// sorting part of the arrays or you will -/// have bugs! -/// @return void -/////////////////////////////////////////////////////////////////////////// -template -CARE_INLINE -std::enable_if_t::raw_type>::value, void> -sortKeyValueArrays(host_device_ptr & keys, - host_device_ptr & values, - const size_t start, const size_t len, - const bool noCopy) -{ - bool _noCopy ; - if (noCopy && start > 0) { - printf("[CARE] Warning: sortKeyValueArrays. noCopy should not be set if start > 0 (%d)\n", (int)start); - _noCopy = false; - } - else { - _noCopy = noCopy; - } - - if constexpr (std::is_same_v) { - detail::stableSortKeyValuePairs(keys, values, len, start); - } - else { - // TODO openMP parallel implementation -#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))) - - // Allocate space for the result - host_device_ptr keyResult{len}; - host_device_ptr valueResult{len}; - - // Get the raw data to pass to cub - CHAIDataGetter keyGetter {}; - CHAIDataGetter valueGetter {}; - - auto * rawKeyData = keyGetter.getRawArrayData(keys) + start; - auto * rawValueData = valueGetter.getRawArrayData(values) + start; - - auto * rawKeyResult = keyGetter.getRawArrayData(keyResult); - auto * rawValueResult = valueGetter.getRawArrayData(valueResult); - - using RawKeyType = std::remove_reference_t; - - auto custom_comparator = [] CARE_HOST_DEVICE (const RawKeyType& lhs, - const RawKeyType& rhs) { - return lhs < rhs; - }; - - // Get the temp storage length - char * d_temp_storage = nullptr; - size_t temp_storage_bytes = 0; - - // When called with a nullptr for temp storage, this returns how much - // temp storage should be allocated. - if (len > 0) { -#if defined(__CUDACC__) - cub::DeviceMergeSort::StableSortPairs((void *)d_temp_storage, temp_storage_bytes, - rawKeyData, - rawValueData, - len, custom_comparator); -#elif defined(__HIPCC__) - hipcub::DeviceMergeSort::StableSortPairs((void *)d_temp_storage, temp_storage_bytes, - rawKeyData, - rawValueData, - len, custom_comparator); -#endif - } - - // Allocate the temp storage and get raw data to pass to cub - host_device_ptr tmpManaged {temp_storage_bytes}; - - CHAIDataGetter charGetter {}; - d_temp_storage = charGetter.getRawArrayData(tmpManaged); - - // Now sort - if (len > 0) { -#if defined(CHAI_THIN_GPU_ALLOCATE) - chai::ArrayManager::getInstance()->setExecutionSpace(chai::GPU); -#endif - -#if defined(__CUDACC__) - cub::DeviceMergeSort::StableSortPairs((void *)d_temp_storage, temp_storage_bytes, - rawKeyData, - rawValueData, - len, - custom_comparator); -#elif defined(__HIPCC__) - hipcub::DeviceMergeSort::StableSortPairs((void *)d_temp_storage, temp_storage_bytes, - rawKeyData, - rawValueData, - len, - custom_comparator); -#endif - -#if defined(CHAI_THIN_GPU_ALLOCATE) - chai::ArrayManager::getInstance()->setExecutionSpace(chai::NONE); -#endif - - tmpManaged.free(); - } - - // merge sort did an inplace sort, so the answer is already in keys and Values - if (len > 0) { - keyResult.free(); - valueResult.free(); - } - -#else // defined(CARE_GPUCC) - // fall back to an implementation that uses std::stable_sort on the host - detail::stableSortKeyValuePairs(keys, values, len, start); - -#endif // defined(CARE_GPUCC) - } -} - - #if defined(CARE_PARALLEL_DEVICE) || CARE_ENABLE_GPU_SIMULATION_MODE /////////////////////////////////////////////////////////////////////////// /// @author Benjamin Liu after Alan Dayton /// @brief Initializes keys and values by copying elements from the array /// @param[out] keys - The key array to set to the identity /// @param[out] values - The value array to set -/// @param[in] len - The number of elements to allocate space for -/// @param[in] arr - An array to copy elements from +/// @param[in] len - The number of elements to copy +/// @param[in] arr - input array /// @return void /////////////////////////////////////////////////////////////////////////// template @@ -356,7 +68,7 @@ CARE_INLINE void setKeyValueArraysFromArray(host_device_ptr & keys, /// @brief Initializes the KeyValueSorter by copying elements from the array /// @param[out] keys - The key array to set to the identity /// @param[out] values - The value array to set -/// @param[in] len - The number of elements to allocate space for +/// @param[in] len - The number of elements to copy /// @param[in] arr - An array to copy elements from /// @return void /////////////////////////////////////////////////////////////////////////// @@ -376,16 +88,13 @@ CARE_INLINE void setKeyValueArraysFromManagedArray(host_device_ptr & ke /////////////////////////////////////////////////////////////////////////// /// @author Jeff Keasler, Alan Dayton -/// @brief Eliminates duplicate values -/// Remove duplicate values from old key/value arrays. -/// Old key/value arrays should already be sorted by value. -/// New key/value arrays should be allocated to the old size. +/// @brief Eliminates duplicate values from sorted key/value arrays /// @param[out] newKeys New key array with duplicates removed /// @param[out] newValues New value array with duplicates removed -/// @param[in] oldKeys Old key array (key-value pairs sorted by value) +/// @param[in] oldKeys Old key array /// @param[in] oldValues Old value array (sorted) -/// @param[in] oldLen Length of old key/value array and initial length for new -/// @return newLen Length of new key/value arrays +/// @param[in] oldLen Length of the old arrays +/// @return Length of the new key/value arrays /////////////////////////////////////////////////////////////////////////// template CARE_INLINE size_t eliminateKeyValueDuplicates(host_device_ptr& newKeys, @@ -409,34 +118,33 @@ CARE_INLINE size_t eliminateKeyValueDuplicates(host_device_ptr& newKeys return (size_t)newSize; } -template +template CARE_INLINE void IntersectKeyValueSorters(RAJADeviceExec exec, - KeyValueSorter sorter1, int size1, - KeyValueSorter sorter2, int size2, + KeyValueSorter sorter1, SizeType size1, + KeyValueSorter sorter2, SizeType size2, host_device_ptr& matches1, host_device_ptr& matches2, - int & numMatches) + SizeType & numMatches) { - int smaller = (size1 < size2) ? size1 : size2 ; - int start1 = 0; - int start2 = 0; + SizeType smaller = (size1 < size2) ? size1 : size2; + SizeType start1 = 0; + SizeType start2 = 0; - numMatches = 0 ; + numMatches = 0; if (smaller == 0) { - matches1 = nullptr ; - matches2 = nullptr ; - return ; - } - else { - matches1.alloc(smaller); - matches1.namePointer("matches1"); - matches2.alloc(smaller); - matches2.namePointer("matches2"); + matches1 = nullptr; + matches2 = nullptr; + return; } + matches1.alloc(smaller); + matches1.namePointer("matches1"); + matches2.alloc(smaller); + matches2.namePointer("matches2"); + host_device_ptr smallerMatches, largerMatches; host_device_ptr smallerKeys, largerKeys; - int larger, smallStart, largeStart; + SizeType larger, smallStart, largeStart; host_device_ptr smallerArray, largerArray; if (smaller == size1) { smallerArray = sorter1.values(); @@ -461,16 +169,16 @@ CARE_INLINE void IntersectKeyValueSorters(RAJADeviceExec exec, largerMatches = matches1; } - host_device_ptr searches(smaller+1); - host_device_ptr matched(smaller+1); - CARE_STREAM_LOOP(i, 0, smaller+1) { + host_device_ptr searches(smaller + 1); + host_device_ptr matched(smaller + 1); + CARE_STREAM_LOOP(i, 0, smaller + 1) { if (i == smaller) { searches[i] = -1; } else { // to be consistent with CPU algorithm, find the first match - int match = care::BinarySearch(largerArray, largeStart, larger, smallerArray[i+smallStart]); - while (match > largeStart && largerArray[match-1] == largerArray[match]) { + int match = care::BinarySearch(largerArray, largeStart, larger, smallerArray[i + smallStart]); + while (match > largeStart && largerArray[match - 1] == largerArray[match]) { --match; } searches[i] = match; @@ -478,28 +186,27 @@ CARE_INLINE void IntersectKeyValueSorters(RAJADeviceExec exec, matched[i] = i != smaller && searches[i] > -1; } CARE_STREAM_LOOP_END - care::exclusive_scan(RAJADeviceExec{}, matched, nullptr, smaller+1, 0, true); + care::exclusive_scan(RAJADeviceExec{}, matched, nullptr, smaller + 1, 0, true); CARE_STREAM_LOOP(i, 0, smaller) { if (searches[i] > -1) { - smallerMatches[matched[i]] = smallerKeys[i+smallStart]; + smallerMatches[matched[i]] = smallerKeys[i + smallStart]; largerMatches[matched[i]] = largerKeys[searches[i]]; } } CARE_STREAM_LOOP_END - numMatches = matched.pick(smaller); + numMatches = matched.pick(smaller); searches.free(); matched.free(); - /* change the size of the array */ if (numMatches == 0) { matches1.free(); matches2.free(); } else { + /* reduce the size of the matches arrays*/ matches1.realloc(numMatches); matches2.realloc(numMatches); } - } #endif // defined(CARE_PARALLEL_DEVICE) || CARE_ENABLE_GPU_SIMULATION_MODE diff --git a/src/care/LoopFuser.cpp b/src/care/LoopFuser.cpp index 88c39f4e..d0808900 100644 --- a/src/care/LoopFuser.cpp +++ b/src/care/LoopFuser.cpp @@ -283,14 +283,15 @@ void LoopFuser::flush_parallel_scans(const char * fileN /* need to write the scan positions to the output destinations */ /* each destination is computed */ + care::host_ptr * pos_output_destinations = m_pos_output_destinations; CARE_SEQUENTIAL_LOOP(actionIndex, 0, action_count) { int scan_pos_offset = actionIndex == 0 ? 0 : scan_pos_outputs[actionIndex-1]; int pos = scan_pos_outputs[actionIndex]; pos -= scan_pos_offset; - *(m_pos_output_destinations[actionIndex].data()) += pos; + *(pos_output_destinations[actionIndex].data()) += pos; if (very_verbose) { printf("actionIndex %i: scan_pos_offset %i scan_pos_output %i pos %i store %i \n", - actionIndex, scan_pos_offset, scan_pos_outputs[actionIndex], pos, *(m_pos_output_destinations[actionIndex].data())); + actionIndex, scan_pos_offset, scan_pos_outputs[actionIndex], pos, *(pos_output_destinations[actionIndex].data())); } } CARE_SEQUENTIAL_LOOP_END scan_var.free(); diff --git a/src/care/algorithm_decl.h b/src/care/algorithm_decl.h index a84959e4..48a48b1a 100644 --- a/src/care/algorithm_decl.h +++ b/src/care/algorithm_decl.h @@ -493,6 +493,310 @@ template void ExpandArrayInPlace(RAJADeviceExec, care::host_device_ptr array, care::host_device_ptr indexSet, int length); #endif // defined(CARE_PARALLEL_DEVICE) + +/////////////////////////////////////////////////////////////////////////// +/// @author Ben Liu, Peter Robinson, Alan Dayton +/// @brief Checks whether an array of type T is sorted and optionally unique. +/// @param[in] array - The array to check +/// @param[in] len - The number of elements contained in the sorter +/// @param[in] name - The name of the calling function +/// @param[in] argname - The name of the sorter in the calling function +/// @param[in] allowDuplicates - Whether or not to allow duplicates +/// @param[in] warnOnFailure - Whether to print a warning if array not sorted +/// @return true if sorted, false otherwise +/////////////////////////////////////////////////////////////////////////// +template +CARE_HOST_DEVICE CARE_INLINE bool checkSorted(const T* array, const int len, + const char* name, const char* argname, + const bool allowDuplicates, + const bool warnOnFailure) +{ + if (len > 0) { + int last = 0; + bool failed = false; + + if (allowDuplicates) { + for (int k = 1 ; k < len ; ++k) { + failed = array[k] < array[last]; + + if (failed) { + break; + } + else { + last = k; + } + } + } + else { + for (int k = 1 ; k < len ; ++k) { + failed = array[k] <= array[last]; + + if (failed) { + break; + } + else { + last = k; + } + } + } + + if (failed) { + if (warnOnFailure) { + printf("care:%s: %s not in ascending order at index %d\n", name, argname, last + 1); + } + return false; + } + } + + return true; +} + +template +CARE_HOST_DEVICE CARE_INLINE bool checkSorted(const care::host_device_ptr& array, + const int len, + const char* name, + const char* argname, + const bool allowDuplicates, + const bool warnOnFailure) +{ + return checkSorted(array.data(), len, name, argname, allowDuplicates, warnOnFailure); +} + +/************************************************************************ + * Function : BinarySearch + * Author(s) : Brad Wallin, Peter Robinson + * Purpose : Every good code has to have one. Searches a sorted array, + * or a sorted subarray, for a particular value. This used to + * be in NodesGlobalToLocal. The algorithm was taken from + * Numerical Recipes in C, Second Edition. + * + * Important Note: mapSize is the length of the region you + * are searching. For example, if you have an array that has + * 100 entries in it, and you want to search from index 5 to + * 40, then you would set start=5, and mapSize=(40-5)=35. + * In other words, mapSize is NOT the original length of the + * array and it is also NOT the ending index for your search. + * + * If returnUpperBound is set to true, this will return the + * index corresponding to the earliest entry that is greater + * than num. A return value of -1 indicates that all values + * in map are smaller than or equal to num. + * + * @NOTE: Intentionally implemented this using only the '<' + * operator to follow weak strict ordering semantics. + * + ************************************************************************/ + +template +CARE_HOST_DEVICE CARE_INLINE int BinarySearch(const T *map, const int start, + const int mapSize, const T num, + bool returnUpperBound) +{ + int klo = start ; + int khi = start + mapSize; + int k = ((khi+klo) >> 1) + 1 ; + + if ((map == nullptr) || (mapSize == 0)) { + return -1 ; + } +#ifdef CARE_DEBUG + const bool allowDuplicates = true; + const bool warnOnFailure = true; + checkSorted(&(map[start]), mapSize, "BinarySearch", "map", allowDuplicates, warnOnFailure) ; +#endif + + while (khi-klo > 1) { + k = (khi+klo) >> 1 ; + if (! (map[k] < num) && !(num < map[k])) { + if (returnUpperBound) { + khi = k+1; + klo = k; + continue; + } + else { + return k ; + } + } + else if (num < map[k]) { + khi = k ; + } + else { + klo = k ; + } + } + if (returnUpperBound) { + k = klo; + // the lower option bounds num + if (num < map[k]) { + return k; + } + // the upper option is within the range of the map index set + if (khi < start + mapSize) { + // Note: fix for last test in TEST(algorithm, binarysearch). This algorithm has failed to pick up the upper + // bound above 1 in the array {0, 1, 1, 1, 1, 1, 6}. Having 1 repeated confused the algorithm. + while ((khi < start + mapSize) && (!(map[khi] < num) && !(num < map[khi]))) { + ++khi; + } + + // the upper option bounds num + if ((khi < start + mapSize) && (num < map[khi])) { + return khi; + } + // neither the upper or lower option bound num + return -1; + } + else { + // the lower option does not bound num, and the upper option is out of bounds + return -1; + } + } + --k; + if (!(map[k] < num) && !(num < map[k])) { + return k ; + } + else { + return -1 ; + } +} + +template +CARE_HOST_DEVICE CARE_INLINE int BinarySearch(const care::host_device_ptr& map, const int start, + const int mapSize, const mapType num, + bool returnUpperBound) +{ + return BinarySearch(map.data(), start, mapSize, num, returnUpperBound); +} + +template +CARE_HOST_DEVICE CARE_INLINE int BinarySearch(const care::host_device_ptr& map, const int start, + const int mapSize, const mapType num, + bool returnUpperBound) +{ + return BinarySearch(map.data(), start, mapSize, num, returnUpperBound); +} + + +/************************************************************************ + * Function : uniqLocal + * Author(s) : Benjamin Liu + * Purpose : Remove duplicates in-place from an array that is sorted + * in ascending order and updates len. + * For calls from within RAJA loops. + * Does not reallocate array. + ************************************************************************/ +template +CARE_HOST_DEVICE CARE_INLINE void uniqLocal(care::local_ptr array, int& len) +{ + int origLen = len ; + len = 0 ; + + int i = 0 ; + while (i < origLen) { + /* copy the unique value into the array */ + array[len] = array[i] ; + /* skip over all the redundant elements */ + while (i < origLen && array[i] == array[len]) { + ++i ; + } + ++len ; + } +} + +template +CARE_HOST_DEVICE CARE_INLINE T ArrayMin(care::local_ptr arr, int n, T initVal, int startIndex) +{ + T min = initVal; + for (int k = startIndex; k < n; ++k) { + min = care::min(min, arr[k]); + } + return min; +} + +template +CARE_HOST_DEVICE CARE_INLINE T ArrayMin(care::local_ptr arr, int n, T initVal, int startIndex) +{ + return ArrayMin((care::local_ptr)arr, n, initVal, startIndex); +} + +template +CARE_HOST_DEVICE CARE_INLINE T ArrayMax(care::local_ptr arr, int n, T initVal, int startIndex) +{ + T max = initVal; + for (int k = startIndex; k < n; ++k) { + max = care::max(max, arr[k]); + } + return max; +} + +template +CARE_HOST_DEVICE CARE_INLINE T ArrayMax(care::local_ptr arr, int n, T initVal, int startIndex) +{ + return ArrayMax((care::local_ptr)arr, n, initVal, startIndex); +} + +/************************************************************************ + * Function : ArrayMinMax + * Author(s) : Peter Robinson + * Purpose : Stores Minimum / Maximum values of arr (as a double) in outMin / outMax; + * If mask was such that no values were compared, returns 0, outMin will be -DBL_MAX, outMax will be DBL_MAX + * Otherwise, returns 1. + * care::local_ptr API to support calls from within RAJA contexts. + * ************************************************************************/ +template +CARE_HOST_DEVICE CARE_INLINE int ArrayMinMax(care::local_ptr arr, + care::local_ptr mask, + int n, double *outMin, double *outMax) +{ + bool result = false; + // a previous implementation had min and max as a templated type and then used std::numeric_limits::lowest() and + // std::numeric_limits::max() for initial values, but that is not valid on the device and results in + // warnings and undefined behavior at runtime. + double min, max; + if (arr) { + max = -DBL_MAX; + min = DBL_MAX; + if (mask) { + for (int i = 0; i < n; ++i) { + if (mask[i]) { + min = care::min(min, (double)arr[i]); + max = care::max(max, (double)arr[i]); + } + } + if (min != DBL_MAX || + max != -DBL_MAX) { + result = true; + } + } + else { + for (int i = 0; i < n; ++i) { + min = care::min(min, (double)arr[i]); + max = care::max(max, (double)arr[i]); + } + result = true; + } + } + + if (result) { + *outMin = (double) min; + *outMax = (double) max; + } + else { + *outMin = -DBL_MAX; + *outMax = +DBL_MAX; + } + return (int) result; +} + +template +CARE_HOST_DEVICE CARE_INLINE int ArrayMinMax(care::local_ptr arr, + care::local_ptr mask, + int n, double *outMin, double *outMax) +{ + return ArrayMinMax((care::local_ptr)arr, (care::local_ptr)mask, n, outMin, outMax); +} + + + } // end namespace care #endif // !defined(CARE_ALGORITHM_DECL_H) diff --git a/src/care/algorithm_impl.h b/src/care/algorithm_impl.h index 15226388..67540dea 100644 --- a/src/care/algorithm_impl.h +++ b/src/care/algorithm_impl.h @@ -36,73 +36,6 @@ namespace care { -/////////////////////////////////////////////////////////////////////////// -/// @author Ben Liu, Peter Robinson, Alan Dayton -/// @brief Checks whether an array of type T is sorted and optionally unique. -/// @param[in] array - The array to check -/// @param[in] len - The number of elements contained in the sorter -/// @param[in] name - The name of the calling function -/// @param[in] argname - The name of the sorter in the calling function -/// @param[in] allowDuplicates - Whether or not to allow duplicates -/// @param[in] warnOnFailure - Whether to print a warning if array not sorted -/// @return true if sorted, false otherwise -/////////////////////////////////////////////////////////////////////////// -template -CARE_HOST_DEVICE CARE_INLINE bool checkSorted(const T* array, const int len, - const char* name, const char* argname, - const bool allowDuplicates, - const bool warnOnFailure) -{ - if (len > 0) { - int last = 0; - bool failed = false; - - if (allowDuplicates) { - for (int k = 1 ; k < len ; ++k) { - failed = array[k] < array[last]; - - if (failed) { - break; - } - else { - last = k; - } - } - } - else { - for (int k = 1 ; k < len ; ++k) { - failed = array[k] <= array[last]; - - if (failed) { - break; - } - else { - last = k; - } - } - } - - if (failed) { - if (warnOnFailure) { - printf("care:%s: %s not in ascending order at index %d\n", name, argname, last + 1); - } - return false; - } - } - - return true; -} - -template -CARE_HOST_DEVICE CARE_INLINE bool checkSorted(const care::host_device_ptr& array, - const int len, - const char* name, - const char* argname, - const bool allowDuplicates, - const bool warnOnFailure) -{ - return checkSorted(array.data(), len, name, argname, allowDuplicates, warnOnFailure); -} /************************************************************************ * Function : IntersectArrays @@ -417,118 +350,6 @@ CARE_INLINE void IntersectArrays(RAJA::seq_exec exec, matches1, matches2, numMatches); } -/************************************************************************ - * Function : BinarySearch - * Author(s) : Brad Wallin, Peter Robinson - * Purpose : Every good code has to have one. Searches a sorted array, - * or a sorted subarray, for a particular value. This used to - * be in NodesGlobalToLocal. The algorithm was taken from - * Numerical Recipes in C, Second Edition. - * - * Important Note: mapSize is the length of the region you - * are searching. For example, if you have an array that has - * 100 entries in it, and you want to search from index 5 to - * 40, then you would set start=5, and mapSize=(40-5)=35. - * In other words, mapSize is NOT the original length of the - * array and it is also NOT the ending index for your search. - * - * If returnUpperBound is set to true, this will return the - * index corresponding to the earliest entry that is greater - * than num. A return value of -1 indicates that all values - * in map are smaller than or equal to num. - * - * @NOTE: Intentionally implemented this using only the '<' - * operator to follow weak strict ordering semantics. - * - ************************************************************************/ - -template -CARE_HOST_DEVICE CARE_INLINE int BinarySearch(const T *map, const int start, - const int mapSize, const T num, - bool returnUpperBound) -{ - int klo = start ; - int khi = start + mapSize; - int k = ((khi+klo) >> 1) + 1 ; - - if ((map == nullptr) || (mapSize == 0)) { - return -1 ; - } -#ifdef CARE_DEBUG - const bool allowDuplicates = true; - const bool warnOnFailure = true; - checkSorted(&(map[start]), mapSize, "BinarySearch", "map", allowDuplicates, warnOnFailure) ; -#endif - - while (khi-klo > 1) { - k = (khi+klo) >> 1 ; - if (! (map[k] < num) && !(num < map[k])) { - if (returnUpperBound) { - khi = k+1; - klo = k; - continue; - } - else { - return k ; - } - } - else if (num < map[k]) { - khi = k ; - } - else { - klo = k ; - } - } - if (returnUpperBound) { - k = klo; - // the lower option bounds num - if (num < map[k]) { - return k; - } - // the upper option is within the range of the map index set - if (khi < start + mapSize) { - // Note: fix for last test in TEST(algorithm, binarysearch). This algorithm has failed to pick up the upper - // bound above 1 in the array {0, 1, 1, 1, 1, 1, 6}. Having 1 repeated confused the algorithm. - while ((khi < start + mapSize) && (!(map[khi] < num) && !(num < map[khi]))) { - ++khi; - } - - // the upper option bounds num - if ((khi < start + mapSize) && (num < map[khi])) { - return khi; - } - // neither the upper or lower option bound num - return -1; - } - else { - // the lower option does not bound num, and the upper option is out of bounds - return -1; - } - } - --k; - if (!(map[k] < num) && !(num < map[k])) { - return k ; - } - else { - return -1 ; - } -} - -template -CARE_HOST_DEVICE CARE_INLINE int BinarySearch(const care::host_device_ptr& map, const int start, - const int mapSize, const mapType num, - bool returnUpperBound) -{ - return BinarySearch(map.data(), start, mapSize, num, returnUpperBound); -} - -template -CARE_HOST_DEVICE CARE_INLINE int BinarySearch(const care::host_device_ptr& map, const int start, - const int mapSize, const mapType num, - bool returnUpperBound) -{ - return BinarySearch(map.data(), start, mapSize, num, returnUpperBound); -} #ifdef CARE_PARALLEL_DEVICE /************************************************************************ @@ -863,31 +684,6 @@ CARE_INLINE void sort_uniq(Exec e, care::host_device_ptr * array, int * len, *len = uniqArray(e, *array, *len, noCopy); } -/************************************************************************ - * Function : uniqLocal - * Author(s) : Benjamin Liu - * Purpose : Remove duplicates in-place from an array that is sorted - * in ascending order and updates len. - * For calls from within RAJA loops. - * Does not reallocate array. - ************************************************************************/ -template -CARE_HOST_DEVICE CARE_INLINE void uniqLocal(care::local_ptr array, int& len) -{ - int origLen = len ; - len = 0 ; - - int i = 0 ; - while (i < origLen) { - /* copy the unique value into the array */ - array[len] = array[i] ; - /* skip over all the redundant elements */ - while (i < origLen && array[i] == array[len]) { - ++i ; - } - ++len ; - } -} template CARE_INLINE void ExpandArrayInPlace(RAJA::seq_exec, care::host_device_ptr array, @@ -981,21 +777,6 @@ CARE_INLINE T ArrayMin(care::host_device_ptr arr, int n, T initVal, int start return ArrayMin((care::host_device_ptr)arr, n, initVal, startIndex); } -template -CARE_HOST_DEVICE CARE_INLINE T ArrayMin(care::local_ptr arr, int n, T initVal, int startIndex) -{ - T min = initVal; - for (int k = startIndex; k < n; ++k) { - min = care::min(min, arr[k]); - } - return min; -} - -template -CARE_HOST_DEVICE CARE_INLINE T ArrayMin(care::local_ptr arr, int n, T initVal, int startIndex) -{ - return ArrayMin((care::local_ptr)arr, n, initVal, startIndex); -} /************************************************************************ * Function : ArrayMin @@ -1052,22 +833,6 @@ CARE_INLINE T ArrayMax(care::host_device_ptr arr, int n, T initVal, int start return ArrayMax((care::host_device_ptr)arr, n, initVal, startIndex); } -template -CARE_HOST_DEVICE CARE_INLINE T ArrayMax(care::local_ptr arr, int n, T initVal, int startIndex) -{ - T max = initVal; - for (int k = startIndex; k < n; ++k) { - max = care::max(max, arr[k]); - } - return max; -} - -template -CARE_HOST_DEVICE CARE_INLINE T ArrayMax(care::local_ptr arr, int n, T initVal, int startIndex) -{ - return ArrayMax((care::local_ptr)arr, n, initVal, startIndex); -} - /************************************************************************ * Function : ArrayMax * Author(s) : Peter Robinson @@ -1202,67 +967,6 @@ CARE_INLINE int ArrayMinMax(care::host_device_ptr arr, #endif // CARE_HAVE_LLNL_GLOBALID -/************************************************************************ - * Function : ArrayMinMax - * Author(s) : Peter Robinson - * Purpose : Stores Minimum / Maximum values of arr (as a double) in outMin / outMax; - * If mask was such that no values were compared, returns 0, outMin will be -DBL_MAX, outMax will be DBL_MAX - * Otherwise, returns 1. - * care::local_ptr API to support calls from within RAJA contexts. - * ************************************************************************/ -template -CARE_HOST_DEVICE CARE_INLINE int ArrayMinMax(care::local_ptr arr, - care::local_ptr mask, - int n, double *outMin, double *outMax) -{ - bool result = false; - // a previous implementation had min and max as a templated type and then used std::numeric_limits::lowest() and - // std::numeric_limits::max() for initial values, but that is not valid on the device and results in - // warnings and undefined behavior at runtime. - double min, max; - if (arr) { - max = -DBL_MAX; - min = DBL_MAX; - if (mask) { - for (int i = 0; i < n; ++i) { - if (mask[i]) { - min = care::min(min, (double)arr[i]); - max = care::max(max, (double)arr[i]); - } - } - if (min != DBL_MAX || - max != -DBL_MAX) { - result = true; - } - } - else { - for (int i = 0; i < n; ++i) { - min = care::min(min, (double)arr[i]); - max = care::max(max, (double)arr[i]); - } - result = true; - } - } - - if (result) { - *outMin = (double) min; - *outMax = (double) max; - } - else { - *outMin = -DBL_MAX; - *outMax = +DBL_MAX; - } - return (int) result; -} - -template -CARE_HOST_DEVICE CARE_INLINE int ArrayMinMax(care::local_ptr arr, - care::local_ptr mask, - int n, double *outMin, double *outMax) -{ - return ArrayMinMax((care::local_ptr)arr, (care::local_ptr)mask, n, outMin, outMax); -} - /************************************************************************ * Function : ArrayCount * Author(s) : Peter Robinson diff --git a/src/care/care_inst.h b/src/care/care_inst.h index 0f7511ef..ef4c5d04 100644 --- a/src/care/care_inst.h +++ b/src/care/care_inst.h @@ -116,30 +116,6 @@ namespace care { /////////////////////////////////////////////////////////////////////////////// -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE bool checkSorted(const int*, const int, const char*, const char*, const bool, const bool) ; -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE bool checkSorted(const float*, const int, const char*, const char*, const bool, const bool) ; -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE bool checkSorted(const double*, const int, const char*, const char*, const bool, const bool) ; -#if CARE_HAVE_LLNL_GLOBALID -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE bool checkSorted(const globalID*, const int, const char*, const char*, const bool, const bool) ; -#endif - -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE bool checkSorted(const care::host_device_ptr&, const int, const char*, const char*, const bool, const bool) ; -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE bool checkSorted(const care::host_device_ptr&, const int, const char*, const char*, const bool, const bool) ; -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE bool checkSorted(const care::host_device_ptr&, const int, const char*, const char*, const bool, const bool) ; -#if CARE_HAVE_LLNL_GLOBALID -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE bool checkSorted(const care::host_device_ptr&, const int, const char*, const char*, const bool, const bool) ; -#endif - -/////////////////////////////////////////////////////////////////////////////// - #ifdef CARE_PARALLEL_DEVICE CARE_EXTERN template CARE_DLL_API @@ -188,53 +164,6 @@ void IntersectArrays(RAJA::seq_exec, care::host_device_ptr, int, int, /////////////////////////////////////////////////////////////////////////////// -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE int BinarySearch(const int *, const int, const int, const int, bool) ; -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE int BinarySearch(const size_t *, const int, const int, const size_t, bool) ; -#if CARE_HAVE_LLNL_GLOBALID - -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE int BinarySearch(const globalID *, const int, const int, const globalID, bool) ; -#if GLOBALID_IS_64BIT -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE int BinarySearch(const GIDTYPE *, const int, const int, const GIDTYPE, bool) ; -#endif - -#endif - -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE int BinarySearch(const care::host_device_ptr&, const int, const int, const int, bool) ; -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE int BinarySearch(const care::host_device_ptr&, const int, const int, const size_t, bool) ; -#if CARE_HAVE_LLNL_GLOBALID - -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE int BinarySearch(const care::host_device_ptr&, const int, const int, const globalID, bool) ; -#if GLOBALID_IS_64BIT -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE int BinarySearch(const care::host_device_ptr&, const int, const int, const GIDTYPE, bool) ; -#endif - -#endif - -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE int BinarySearch(const care::host_device_ptr&, const int, const int, const int, bool) ; -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE int BinarySearch(const care::host_device_ptr&, const int, const int, const size_t, bool) ; -#if CARE_HAVE_LLNL_GLOBALID - -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE int BinarySearch(const care::host_device_ptr&, const int, const int, const globalID, bool) ; -#if GLOBALID_IS_64BIT -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE int BinarySearch(const care::host_device_ptr&, const int, const int, const GIDTYPE, bool) ; -#endif - -#endif - -/////////////////////////////////////////////////////////////////////////////// - #ifdef CARE_PARALLEL_DEVICE CARE_EXTERN template CARE_DLL_API @@ -421,31 +350,6 @@ int CompressArray(care::host_device_ptr &, const int, care::host_devic /////////////////////////////////////////////////////////////////////////////// -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE void InsertionSort(care::local_ptr, int) ; -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE void InsertionSort(care::local_ptr, int) ; -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE void InsertionSort(care::local_ptr, int) ; -#if CARE_HAVE_LLNL_GLOBALID -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE void InsertionSort(care::local_ptr, int) ; -#endif - -/////////////////////////////////////////////////////////////////////////////// - -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE void uniqLocal(care::local_ptr, int&) ; -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE void uniqLocal(care::local_ptr, int&) ; -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE void uniqLocal(care::local_ptr, int&) ; -#if CARE_HAVE_LLNL_GLOBALID -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE void uniqLocal(care::local_ptr, int&) ; -#endif - -/////////////////////////////////////////////////////////////////////////////// #ifdef CARE_PARALLEL_DEVICE @@ -619,33 +523,6 @@ CARE_EXTERN template CARE_DLL_API float ArrayMin(care::host_device_ptr, int, float, int) ; CARE_EXTERN template CARE_DLL_API double ArrayMin(care::host_device_ptr, int, double, int) ; -// TODO GID not implemented - -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE bool ArrayMin(care::local_ptr, int, bool, int) ; -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE int ArrayMin(care::local_ptr, int, int, int) ; -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE float ArrayMin(care::local_ptr, int, float, int) ; -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE double ArrayMin(care::local_ptr, int, double, int) ; -#if CARE_HAVE_LLNL_GLOBALID -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE globalID ArrayMin(care::local_ptr, int, globalID, int) ; -#endif - -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE bool ArrayMin(care::local_ptr, int, bool, int) ; -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE int ArrayMin(care::local_ptr, int, int, int) ; -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE float ArrayMin(care::local_ptr, int, float, int) ; -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE double ArrayMin(care::local_ptr, int, double, int) ; -#if CARE_HAVE_LLNL_GLOBALID -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE globalID ArrayMin(care::local_ptr, int, globalID, int) ; -#endif /////////////////////////////////////////////////////////////////////////////// @@ -745,32 +622,6 @@ CARE_EXTERN template CARE_DLL_API double ArrayMax(care::host_device_ptr, int, double, int) ; // TODO GID not implemented -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE bool ArrayMax(care::local_ptr, int, bool, int) ; -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE int ArrayMax(care::local_ptr, int, int, int) ; -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE float ArrayMax(care::local_ptr, int, float, int) ; -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE double ArrayMax(care::local_ptr, int, double, int) ; -#if CARE_HAVE_LLNL_GLOBALID -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE globalID ArrayMax(care::local_ptr, int, globalID, int) ; -#endif - -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE bool ArrayMax(care::local_ptr, int, bool, int) ; -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE int ArrayMax(care::local_ptr, int, int, int) ; -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE float ArrayMax(care::local_ptr, int, float, int) ; -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE double ArrayMax(care::local_ptr, int, double, int) ; -#if CARE_HAVE_LLNL_GLOBALID -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE globalID ArrayMax(care::local_ptr, int, globalID, int) ; -#endif - CARE_EXTERN template CARE_DLL_API bool ArrayMax(care::host_ptr, int, bool, int) ; CARE_EXTERN template CARE_DLL_API @@ -850,28 +701,6 @@ CARE_EXTERN template CARE_DLL_API int ArrayMinMax(care::host_device_ptr, care::host_device_ptr, int, double *, double *) ; #endif -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE int ArrayMinMax(care::local_ptr, care::local_ptr, int, double *, double *) ; -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE int ArrayMinMax(care::local_ptr, care::local_ptr, int, double *, double *) ; -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE int ArrayMinMax(care::local_ptr, care::local_ptr, int, double *, double *) ; -#if CARE_HAVE_LLNL_GLOBALID -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE int ArrayMinMax(care::local_ptr, care::local_ptr, int, double *, double *) ; -#endif - -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE int ArrayMinMax(care::local_ptr, care::local_ptr, int, double *, double *) ; -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE int ArrayMinMax(care::local_ptr, care::local_ptr, int, double *, double *) ; -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE int ArrayMinMax(care::local_ptr, care::local_ptr, int, double *, double *) ; -#if CARE_HAVE_LLNL_GLOBALID -CARE_EXTERN template CARE_DLL_API -CARE_HOST_DEVICE int ArrayMinMax(care::local_ptr, care::local_ptr, int, double *, double *) ; -#endif - /////////////////////////////////////////////////////////////////////////////// #ifdef CARE_PARALLEL_DEVICE