ICODE-ADC/lammps
4
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Update Kokkos library in LAMMPS to v4.1.0

Browse Source
This commit is contained in:
Stan Gerald Moore
2023-06-29 10:42:42 -06:00
parent 170173a213
commit 330107b77b
480 changed files with 24051 additions and 23393 deletions
Download Patch File Download Diff File Expand all files Collapse all files

8
lib/kokkos/algorithms/CMakeLists.txt
View File

@ -1,7 +1,3 @@
KOKKOS_SUBPACKAGE(Algorithms)
IF (NOT Kokkos_INSTALL_TESTING)
ADD_SUBDIRECTORY(src)
ENDIF()
@ -9,7 +5,3 @@ ENDIF()
IF(NOT ((KOKKOS_ENABLE_OPENMPTARGET OR KOKKOS_ENABLE_OPENACC) AND KOKKOS_CXX_COMPILER_ID STREQUAL NVHPC))
KOKKOS_ADD_TEST_DIRECTORIES(unit_tests)
ENDIF()
KOKKOS_SUBPACKAGE_POSTPROCESS()

5
lib/kokkos/algorithms/cmake/Dependencies.cmake
View File

@ -1,5 +0,0 @@
TRIBITS_PACKAGE_DEFINE_DEPENDENCIES(
LIB_REQUIRED_PACKAGES KokkosCore KokkosContainers
LIB_OPTIONAL_TPLS Pthread CUDA HWLOC
TEST_OPTIONAL_TPLS CUSPARSE
)

4
lib/kokkos/algorithms/cmake/KokkosAlgorithms_config.h.in
View File

@ -1,4 +0,0 @@
#ifndef KOKKOS_ALGORITHMS_CONFIG_H
#define KOKKOS_ALGORITHMS_CONFIG_H
#endif

4
lib/kokkos/algorithms/src/CMakeLists.txt
View File

@ -1,6 +1,3 @@
KOKKOS_CONFIGURE_FILE(${PACKAGE_NAME}_config.h)
#I have to leave these here for tribits
KOKKOS_INCLUDE_DIRECTORIES(${CMAKE_CURRENT_BINARY_DIR})
KOKKOS_INCLUDE_DIRECTORIES(${CMAKE_CURRENT_SOURCE_DIR})
@ -9,7 +6,6 @@ KOKKOS_INCLUDE_DIRECTORIES(${CMAKE_CURRENT_SOURCE_DIR})
FILE(GLOB ALGO_HEADERS *.hpp)
FILE(GLOB ALGO_SOURCES *.cpp)
LIST(APPEND ALGO_HEADERS ${CMAKE_CURRENT_BINARY_DIR}/${PACKAGE_NAME}_config.h)
APPEND_GLOB(ALGO_HEADERS ${CMAKE_CURRENT_SOURCE_DIR}/std_algorithms/*.hpp)
APPEND_GLOB(ALGO_HEADERS ${CMAKE_CURRENT_SOURCE_DIR}/std_algorithms/impl/*.hpp)

2
lib/kokkos/algorithms/src/Kokkos_Random.hpp
View File

@ -1514,7 +1514,7 @@ void fill_random(const ExecutionSpace& exec, ViewType a, RandomPool g,
"Kokkos::fill_random",
Kokkos::RangePolicy<ExecutionSpace>(exec, 0, (LDA + 127) / 128),
Impl::fill_random_functor_begin_end<ViewType, RandomPool, 128,
ViewType::Rank, IndexType>(
ViewType::rank, IndexType>(
a, g, begin, end));
}

137
lib/kokkos/algorithms/src/Kokkos_Sort.hpp
View File

@ -66,11 +66,16 @@
#endif
#if defined(KOKKOS_ENABLE_ONEDPL)
#include <oneapi/dpl/execution>
#include <oneapi/dpl/algorithm>
#endif
namespace Kokkos {
namespace Impl {
template <class DstViewType, class SrcViewType, int Rank = DstViewType::Rank>
template <class DstViewType, class SrcViewType, int Rank = DstViewType::rank>
struct CopyOp;
template <class DstViewType, class SrcViewType>
@ -141,8 +146,12 @@ class BinSort {
Kokkos::is_view<SrcViewType>::value,
Kokkos::View<typename SrcViewType::const_data_type,
typename SrcViewType::array_layout,
typename SrcViewType::device_type,
Kokkos::MemoryTraits<Kokkos::RandomAccess> >,
typename SrcViewType::device_type
#if !defined(KOKKOS_COMPILER_NVHPC) // FIXME_NVHPC
,
Kokkos::MemoryTraits<Kokkos::RandomAccess>
#endif
>,
typename SrcViewType::const_type>;
using perm_view_type = typename PermuteViewType::const_type;
@ -221,7 +230,11 @@ class BinSort {
bool sort_within_bins;
public:
BinSort() = default;
#ifdef KOKKOS_ENABLE_DEPRECATED_CODE_4
KOKKOS_DEPRECATED BinSort() = default;
#else
BinSort() = delete;
#endif
//----------------------------------------
// Constructor: takes the keys, the binning_operator and optionally whether to
@ -324,6 +337,10 @@ class BinSort {
template <class ExecutionSpace, class ValuesViewType>
void sort(const ExecutionSpace& exec, ValuesViewType const& values,
int values_range_begin, int values_range_end) const {
if (values.extent(0) == 0) {
return;
}
static_assert(
Kokkos::SpaceAccessibility<ExecutionSpace,
typename Space::memory_space>::accessible,
@ -335,11 +352,6 @@ class BinSort {
"The provided execution space must be able to access the memory space "
"of the View argument!");
using scratch_view_type =
Kokkos::View<typename ValuesViewType::data_type,
typename ValuesViewType::array_layout,
typename ValuesViewType::device_type>;
const size_t len = range_end - range_begin;
const size_t values_len = values_range_end - values_range_begin;
if (len != values_len) {
@ -347,6 +359,9 @@ class BinSort {
"BinSort::sort: values range length != permutation vector length");
}
using scratch_view_type =
Kokkos::View<typename ValuesViewType::data_type,
typename ValuesViewType::device_type>;
scratch_view_type sorted_values(
view_alloc(exec, WithoutInitializing,
"Kokkos::SortImpl::BinSortFunctor::sorted_values"),
@ -451,24 +466,29 @@ class BinSort {
void operator()(const bin_sort_bins_tag& /*tag*/, const int i) const {
auto bin_size = bin_count_const(i);
if (bin_size <= 1) return;
int upper_bound = bin_offsets(i) + bin_size;
bool sorted = false;
while (!sorted) {
sorted = true;
int old_idx = sort_order(bin_offsets(i));
int new_idx = 0;
for (int k = bin_offsets(i) + 1; k < upper_bound; k++) {
new_idx = sort_order(k);
if (!bin_op(keys_rnd, old_idx, new_idx)) {
sort_order(k - 1) = new_idx;
sort_order(k) = old_idx;
sorted = false;
} else {
old_idx = new_idx;
}
constexpr bool use_std_sort =
std::is_same_v<typename exec_space::memory_space, HostSpace>;
int lower_bound = bin_offsets(i);
int upper_bound = lower_bound + bin_size;
// Switching to std::sort for more than 10 elements has been found
// reasonable experimentally.
if (use_std_sort && bin_size > 10) {
if constexpr (use_std_sort) {
std::sort(&sort_order(lower_bound), &sort_order(upper_bound),
[this](int p, int q) { return bin_op(keys_rnd, p, q); });
}
} else {
for (int k = lower_bound + 1; k < upper_bound; ++k) {
int old_idx = sort_order(k);
int j = k - 1;
while (j >= lower_bound) {
int new_idx = sort_order(j);
if (!bin_op(keys_rnd, old_idx, new_idx)) break;
sort_order(j + 1) = new_idx;
--j;
}
sort_order(j + 1) = old_idx;
}
upper_bound--;
}
}
};
@ -481,7 +501,11 @@ struct BinOp1D {
double mul_ = {};
double min_ = {};
BinOp1D() = default;
#ifdef KOKKOS_ENABLE_DEPRECATED_CODE_4
KOKKOS_DEPRECATED BinOp1D() = default;
#else
BinOp1D() = delete;
#endif
// Construct BinOp with number of bins, minimum value and maximum value
BinOp1D(int max_bins__, typename KeyViewType::const_value_type min,
@ -525,7 +549,11 @@ struct BinOp3D {
double mul_[3] = {};
double min_[3] = {};
BinOp3D() = default;
#ifdef KOKKOS_ENABLE_DEPRECATED_CODE_4
KOKKOS_DEPRECATED BinOp3D() = default;
#else
BinOp3D() = delete;
#endif
BinOp3D(int max_bins__[], typename KeyViewType::const_value_type min[],
typename KeyViewType::const_value_type max[]) {
@ -596,6 +624,10 @@ std::enable_if_t<(Kokkos::is_execution_space<ExecutionSpace>::value) &&
memory_space>::accessible)>
sort(const ExecutionSpace& exec,
const Kokkos::View<DataType, Properties...>& view) {
if (view.extent(0) == 0) {
return;
}
using ViewType = Kokkos::View<DataType, Properties...>;
using CompType = BinOp1D<ViewType>;
@ -634,12 +666,44 @@ sort(const ExecutionSpace& exec,
bin_sort.sort(exec, view);
}
#if defined(KOKKOS_ENABLE_ONEDPL)
template <class DataType, class... Properties>
void sort(const Experimental::SYCL& space,
const Kokkos::View<DataType, Properties...>& view) {
if (view.extent(0) == 0) {
return;
}
using ViewType = Kokkos::View<DataType, Properties...>;
static_assert(SpaceAccessibility<Experimental::SYCL,
typename ViewType::memory_space>::accessible,
"SYCL execution space is not able to access the memory space "
"of the View argument!");
auto queue = space.sycl_queue();
auto policy = oneapi::dpl::execution::make_device_policy(queue);
// Can't use Experimental::begin/end here since the oneDPL then assumes that
// the data is on the host.
static_assert(
ViewType::rank == 1 &&
(std::is_same<typename ViewType::array_layout, LayoutRight>::value ||
std::is_same<typename ViewType::array_layout, LayoutLeft>::value),
"SYCL sort only supports contiguous 1D Views.");
const int n = view.extent(0);
oneapi::dpl::sort(policy, view.data(), view.data() + n);
}
#endif
template <class ExecutionSpace, class DataType, class... Properties>
std::enable_if_t<(Kokkos::is_execution_space<ExecutionSpace>::value) &&
(SpaceAccessibility<
HostSpace, typename Kokkos::View<DataType, Properties...>::
memory_space>::accessible)>
sort(const ExecutionSpace&, const Kokkos::View<DataType, Properties...>& view) {
if (view.extent(0) == 0) {
return;
}
auto first = Experimental::begin(view);
auto last = Experimental::end(view);
std::sort(first, last);
@ -649,6 +713,9 @@ sort(const ExecutionSpace&, const Kokkos::View<DataType, Properties...>& view) {
template <class DataType, class... Properties>
void sort(const Cuda& space,
const Kokkos::View<DataType, Properties...>& view) {
if (view.extent(0) == 0) {
return;
}
const auto exec = thrust::cuda::par.on(space.cuda_stream());
auto first = Experimental::begin(view);
auto last = Experimental::end(view);
@ -659,6 +726,11 @@ void sort(const Cuda& space,
template <class ViewType>
void sort(ViewType const& view) {
Kokkos::fence("Kokkos::sort: before");
if (view.extent(0) == 0) {
return;
}
typename ViewType::execution_space exec;
sort(exec, view);
exec.fence("Kokkos::sort: fence after sorting");
@ -668,6 +740,10 @@ template <class ExecutionSpace, class ViewType>
std::enable_if_t<Kokkos::is_execution_space<ExecutionSpace>::value> sort(
const ExecutionSpace& exec, ViewType view, size_t const begin,
size_t const end) {
if (view.extent(0) == 0) {
return;
}
using range_policy = Kokkos::RangePolicy<typename ViewType::execution_space>;
using CompType = BinOp1D<ViewType>;
@ -690,6 +766,11 @@ std::enable_if_t<Kokkos::is_execution_space<ExecutionSpace>::value> sort(
template <class ViewType>
void sort(ViewType view, size_t const begin, size_t const end) {
Kokkos::fence("Kokkos::sort: before");
if (view.extent(0) == 0) {
return;
}
typename ViewType::execution_space exec;
sort(exec, view, begin, end);
exec.fence("Kokkos::Sort: fence after sorting");

11
lib/kokkos/algorithms/src/std_algorithms/impl/Kokkos_AdjacentFind.hpp
View File

@ -42,12 +42,13 @@ struct StdAdjacentFindFunctor {
const auto& next_value = m_first[i + 1];
const bool are_equal = m_p(my_value, next_value);
auto rv =
are_equal
? red_value_type{i}
: red_value_type{::Kokkos::reduction_identity<IndexType>::min()};
// FIXME_NVHPC using a ternary operator causes problems
red_value_type value = {::Kokkos::reduction_identity<IndexType>::min()};
if (are_equal) {
value.min_loc_true = i;
}
m_reducer.join(red_value, rv);
m_reducer.join(red_value, value);
}
KOKKOS_FUNCTION

2
lib/kokkos/algorithms/src/std_algorithms/impl/Kokkos_Constraints.hpp
View File

@ -29,7 +29,7 @@ struct is_admissible_to_kokkos_std_algorithms : std::false_type {};
template <typename T>
struct is_admissible_to_kokkos_std_algorithms<
T, std::enable_if_t< ::Kokkos::is_view<T>::value && T::rank == 1 &&
T, std::enable_if_t< ::Kokkos::is_view<T>::value && T::rank() == 1 &&
(std::is_same<typename T::traits::array_layout,
Kokkos::LayoutLeft>::value ||
std::is_same<typename T::traits::array_layout,

2
lib/kokkos/algorithms/src/std_algorithms/impl/Kokkos_ExclusiveScan.hpp
View File

@ -96,6 +96,8 @@ struct ExclusiveScanDefaultFunctor {
KOKKOS_FUNCTION
void join(value_type& update, const value_type& input) const {
if (input.is_initial) return;
if (update.is_initial) {
update.val = input.val;
update.is_initial = false;

8
lib/kokkos/algorithms/src/std_algorithms/impl/Kokkos_FindEnd.hpp
View File

@ -59,9 +59,11 @@ struct StdFindEndFunctor {
}
}
const auto rv =
found ? red_value_type{i}
: red_value_type{::Kokkos::reduction_identity<IndexType>::max()};
// FIXME_NVHPC using a ternary operator causes problems
red_value_type rv = {::Kokkos::reduction_identity<IndexType>::max()};
if (found) {
rv.max_loc_true = i;
}
m_reducer.join(red_value, rv);
}

9
lib/kokkos/algorithms/src/std_algorithms/impl/Kokkos_FindFirstOf.hpp
View File

@ -52,10 +52,11 @@ struct StdFindFirstOfFunctor {
}
}
const auto rv =
found ? red_value_type{i}
: red_value_type{::Kokkos::reduction_identity<IndexType>::min()};
// FIXME_NVHPC using a ternary operator causes problems
red_value_type rv = {::Kokkos::reduction_identity<IndexType>::min()};
if (found) {
rv.min_loc_true = i;
}
m_reducer.join(red_value, rv);
}

9
lib/kokkos/algorithms/src/std_algorithms/impl/Kokkos_FindIfOrNot.hpp
View File

@ -44,10 +44,11 @@ struct StdFindIfOrNotFunctor {
// if doing find_if_not, look for when predicate is false
const bool found_condition = is_find_if ? m_p(my_value) : !m_p(my_value);
auto rv =
found_condition
? red_value_type{i}
: red_value_type{::Kokkos::reduction_identity<IndexType>::min()};
// FIXME_NVHPC using a ternary operator causes problems
red_value_type rv = {::Kokkos::reduction_identity<IndexType>::min()};
if (found_condition) {
rv.min_loc_true = i;
}
m_reducer.join(red_value, rv);
}

2
lib/kokkos/algorithms/src/std_algorithms/impl/Kokkos_InclusiveScan.hpp
View File

@ -90,6 +90,8 @@ struct InclusiveScanDefaultFunctor {
KOKKOS_FUNCTION
void join(value_type& update, const value_type& input) const {
if (input.is_initial) return;
if (update.is_initial) {
update.val = input.val;
} else {

8
lib/kokkos/algorithms/src/std_algorithms/impl/Kokkos_IsPartitioned.hpp
View File

@ -43,8 +43,12 @@ struct StdIsPartitionedFunctor {
::Kokkos::reduction_identity<index_type>::min();
constexpr index_type m_red_id_max =
::Kokkos::reduction_identity<index_type>::max();
auto rv = predicate_value ? red_value_type{i, m_red_id_min}
: red_value_type{m_red_id_max, i};
// FIXME_NVHPC using a ternary operator causes problems
red_value_type rv = {m_red_id_max, i};
if (predicate_value) {
rv = {i, m_red_id_min};
}
m_reducer.join(redValue, rv);
}

82
lib/kokkos/algorithms/src/std_algorithms/impl/Kokkos_IsSortedUntil.hpp
View File

@ -28,33 +28,30 @@ namespace Kokkos {
namespace Experimental {
namespace Impl {
template <class IteratorType, class IndicatorViewType, class ComparatorType>
template <class IteratorType, class ComparatorType, class ReducerType>
struct StdIsSortedUntilFunctor {
using index_type = typename IteratorType::difference_type;
using value_type = typename ReducerType::value_type;
IteratorType m_first;
IndicatorViewType m_indicator;
ComparatorType m_comparator;
ReducerType m_reducer;
KOKKOS_FUNCTION
void operator()(const index_type i, int& update, const bool final) const {
void operator()(const index_type i, value_type& reduction_result) const {
const auto& val_i = m_first[i];
const auto& val_ip1 = m_first[i + 1];
if (m_comparator(val_ip1, val_i)) {
++update;
}
if (final) {
m_indicator(i) = update;
m_reducer.join(reduction_result, i);
}
}
KOKKOS_FUNCTION
StdIsSortedUntilFunctor(IteratorType _first1, IndicatorViewType indicator,
ComparatorType comparator)
: m_first(std::move(_first1)),
m_indicator(std::move(indicator)),
m_comparator(std::move(comparator)) {}
StdIsSortedUntilFunctor(IteratorType first, ComparatorType comparator,
ReducerType reducer)
: m_first(std::move(first)),
m_comparator(std::move(comparator)),
m_reducer(std::move(reducer)) {}
};
template <class ExecutionSpace, class IteratorType, class ComparatorType>
@ -73,40 +70,31 @@ IteratorType is_sorted_until_impl(const std::string& label,
}
/*
use scan and a helper "indicator" view
such that we scan the data and fill the indicator with
partial sum that is always 0 unless we find a pair that
breaks the sorting, so in that case the indicator will
have a 1 starting at the location where the sorting breaks.
So finding that 1 means finding the location we want.
*/
Do a par_reduce computing the *min* index that breaks the sorting.
If such an index is found, then the range is sorted until that element.
If no such index is found, then the range is sorted until the end.
*/
using index_type = typename IteratorType::difference_type;
index_type reduction_result;
::Kokkos::Min<index_type> reducer(reduction_result);
::Kokkos::parallel_reduce(
label,
// use num_elements-1 because each index handles i and i+1
RangePolicy<ExecutionSpace>(ex, 0, num_elements - 1),
// use CTAD
StdIsSortedUntilFunctor(first, comp, reducer), reducer);
// aliases
using indicator_value_type = std::size_t;
using indicator_view_type =
::Kokkos::View<indicator_value_type*, ExecutionSpace>;
using functor_type =
StdIsSortedUntilFunctor<IteratorType, indicator_view_type,
ComparatorType>;
// do scan
// use num_elements-1 because each index handles i and i+1
const auto num_elements_minus_one = num_elements - 1;
indicator_view_type indicator("is_sorted_until_indicator_helper",
num_elements_minus_one);
::Kokkos::parallel_scan(
label, RangePolicy<ExecutionSpace>(ex, 0, num_elements_minus_one),
functor_type(first, indicator, std::move(comp)));
// try to find the first sentinel value, which indicates
// where the sorting condition breaks
namespace KE = ::Kokkos::Experimental;
constexpr indicator_value_type sentinel_value = 1;
auto r =
KE::find(ex, KE::cbegin(indicator), KE::cend(indicator), sentinel_value);
const auto shift = r - ::Kokkos::Experimental::cbegin(indicator);
return first + (shift + 1);
/* If the reduction result is equal to the initial value,
it means the range is sorted until the end */
index_type reduction_result_init;
reducer.init(reduction_result_init);
if (reduction_result == reduction_result_init) {
return last;
} else {
/* If such an index is found, then the range is sorted until there and
we need to return an iterator past the element found so do +1 */
return first + (reduction_result + 1);
}
}
template <class ExecutionSpace, class IteratorType>

14
lib/kokkos/algorithms/src/std_algorithms/impl/Kokkos_LexicographicalCompare.hpp
View File

@ -63,12 +63,14 @@ struct StdLexicographicalCompareFunctor {
const auto& my_value1 = m_first1[i];
const auto& my_value2 = m_first2[i];
bool different = m_comparator(my_value1, my_value2) ||
m_comparator(my_value2, my_value1);
auto rv =
different
? red_value_type{i}
: red_value_type{::Kokkos::reduction_identity<IndexType>::min()};
const bool different = m_comparator(my_value1, my_value2) ||
m_comparator(my_value2, my_value1);
// FIXME_NVHPC using a ternary operator causes problems
red_value_type rv = {::Kokkos::reduction_identity<IndexType>::min()};
if (different) {
rv.min_loc_true = i;
}
m_reducer.join(red_value, rv);
}

9
lib/kokkos/algorithms/src/std_algorithms/impl/Kokkos_Mismatch.hpp
View File

@ -42,10 +42,11 @@ struct StdMismatchRedFunctor {
const auto& my_value1 = m_first1[i];
const auto& my_value2 = m_first2[i];
auto rv =
!m_predicate(my_value1, my_value2)
? red_value_type{i}
: red_value_type{::Kokkos::reduction_identity<IndexType>::min()};
// FIXME_NVHPC using a ternary operator causes problems
red_value_type rv = {i};
if (m_predicate(my_value1, my_value2)) {
rv = {::Kokkos::reduction_identity<IndexType>::min()};
}
m_reducer.join(red_value, rv);
}

19
lib/kokkos/algorithms/src/std_algorithms/impl/Kokkos_PartitionCopy.hpp
View File

@ -31,25 +31,6 @@ template <class ValueType>
struct StdPartitionCopyScalar {
ValueType true_count_;
ValueType false_count_;
// Here we implement the copy assignment operators explicitly for consistency
// with how the Scalar structs are implemented inside
// Kokkos_Parallel_Reduce.hpp.
KOKKOS_FUNCTION
void operator=(const StdPartitionCopyScalar& other) {
true_count_ = other.true_count_;
false_count_ = other.false_count_;
}
// this is needed for
// OpenMPTarget/Kokkos_OpenMPTarget_Parallel.hpp:699:21: error: no viable
// overloaded '=' m_returnvalue = 0;
//
KOKKOS_FUNCTION
void operator=(const ValueType value) {
true_count_ = value;
false_count_ = value;
}
};
template <class IndexType, class FirstFrom, class FirstDestTrue,

11
lib/kokkos/algorithms/src/std_algorithms/impl/Kokkos_PartitionPoint.hpp
View File

@ -39,10 +39,13 @@ struct StdPartitionPointFunctor {
KOKKOS_FUNCTION
void operator()(const index_type i, red_value_type& redValue) const {
const auto predicate_value = m_p(m_first[i]);
auto rv =
predicate_value
? red_value_type{::Kokkos::reduction_identity<index_type>::min()}
: red_value_type{i};
// FIXME_NVHPC using a ternary operator causes problems
red_value_type rv = {i};
if (predicate_value) {
rv = {::Kokkos::reduction_identity<index_type>::min()};
}
m_reducer.join(redValue, rv);
}

9
lib/kokkos/algorithms/src/std_algorithms/impl/Kokkos_Reverse.hpp
View File

@ -39,16 +39,7 @@ struct StdReverseFunctor {
KOKKOS_FUNCTION
void operator()(index_type i) const {
// the swap below is doing the same thing, but
// for Intel 18.0.5 does not work.
// But putting the impl directly here, it works.
#ifdef KOKKOS_COMPILER_INTEL
typename InputIterator::value_type tmp = std::move(m_first[i]);
m_first[i] = std::move(m_last[-i - 1]);
m_last[-i - 1] = std::move(tmp);
#else
::Kokkos::Experimental::swap(m_first[i], m_last[-i - 1]);
#endif
}
StdReverseFunctor(InputIterator first, InputIterator last)

8
lib/kokkos/algorithms/src/std_algorithms/impl/Kokkos_Search.hpp
View File

@ -60,9 +60,11 @@ struct StdSearchFunctor {
}
}
const auto rv =
found ? red_value_type{i}
: red_value_type{::Kokkos::reduction_identity<IndexType>::min()};
// FIXME_NVHPC using a ternary operator causes problems
red_value_type rv = {::Kokkos::reduction_identity<IndexType>::min()};
if (found) {
rv = {i};
}
m_reducer.join(red_value, rv);
}

8
lib/kokkos/algorithms/src/std_algorithms/impl/Kokkos_SearchN.hpp
View File

@ -59,9 +59,11 @@ struct StdSearchNFunctor {
}
}
const auto rv =
found ? red_value_type{i}
: red_value_type{::Kokkos::reduction_identity<IndexType>::min()};
// FIXME_NVHPC using a ternary operator causes problems
red_value_type rv = {::Kokkos::reduction_identity<IndexType>::min()};
if (found) {
rv.min_loc_true = i;
}
m_reducer.join(red_value, rv);
}

9
lib/kokkos/algorithms/src/std_algorithms/impl/Kokkos_SwapRanges.hpp
View File

@ -35,16 +35,7 @@ struct StdSwapRangesFunctor {
KOKKOS_FUNCTION
void operator()(IndexType i) const {
// the swap below is doing the same thing, but
// for Intel 18.0.5 does not work.
// But putting the impl directly here, it works.
#ifdef KOKKOS_COMPILER_INTEL
typename IteratorType1::value_type tmp = std::move(m_first1[i]);
m_first1[i] = std::move(m_first2[i]);
m_first2[i] = std::move(tmp);
#else
::Kokkos::Experimental::swap(m_first1[i], m_first2[i]);
#endif
}
KOKKOS_FUNCTION

2
lib/kokkos/algorithms/src/std_algorithms/impl/Kokkos_TransformExclusiveScan.hpp
View File

@ -76,6 +76,8 @@ struct TransformExclusiveScanFunctor {
KOKKOS_FUNCTION
void join(value_type& update, const value_type& input) const {
if (input.is_initial) return;
if (update.is_initial) {
update.val = input.val;
} else {

4
lib/kokkos/algorithms/src/std_algorithms/impl/Kokkos_TransformInclusiveScan.hpp
View File

@ -67,6 +67,8 @@ struct TransformInclusiveScanNoInitValueFunctor {
KOKKOS_FUNCTION
void join(value_type& update, const value_type& input) const {
if (input.is_initial) return;
if (update.is_initial) {
update.val = input.val;
} else {
@ -118,6 +120,8 @@ struct TransformInclusiveScanWithInitValueFunctor {
KOKKOS_FUNCTION
void join(value_type& update, const value_type& input) const {
if (input.is_initial) return;
if (update.is_initial) {
update.val = input.val;
} else {

6
lib/kokkos/algorithms/src/std_algorithms/impl/Kokkos_ValueWrapperForNoNeutralElement.hpp
View File

@ -29,12 +29,6 @@ template <class Scalar>
struct ValueWrapperForNoNeutralElement {
Scalar val;
bool is_initial = true;
KOKKOS_FUNCTION
void operator=(const ValueWrapperForNoNeutralElement& rhs) {
val = rhs.val;
is_initial = rhs.is_initial;
}
};
} // namespace Impl

85
lib/kokkos/algorithms/unit_tests/CMakeLists.txt
View File

@ -16,35 +16,45 @@ foreach(Tag Threads;Serial;OpenMP;Cuda;HPX;HIP;SYCL;OpenMPTarget)
set(dir ${CMAKE_CURRENT_BINARY_DIR}/${dir})
file(MAKE_DIRECTORY ${dir})
# -------------------------
# Sort1d,3d, Random
# -------------------------
set(SOURCES_A)
if(Tag STREQUAL "OpenMP")
LIST(APPEND SOURCES_A
TestOpenMP_Sort1D.cpp
TestOpenMP_Sort3D.cpp
TestOpenMP_SortDynamicView.cpp
)
endif()
# ------------------------------------------
# Sort
# ------------------------------------------
# Each of these inputs is an .hpp file.
# Generate a .cpp file for each one that runs it on the current backend (Tag),
# and add this .cpp file to the sources for UnitTest_RandomAndSort.
foreach(SOURCES_A_Input
TestRandomCommon
TestSortCommon
TestNestedSort
)
set(file ${dir}/${SOURCES_A_Input}.cpp)
set(ALGO_SORT_SOURCES)
foreach(SOURCE_Input
TestSort
TestBinSortA
TestBinSortB
TestNestedSort
)
set(file ${dir}/${SOURCE_Input}.cpp)
# Write to a temporary intermediate file and call configure_file to avoid
# updating timestamps triggering unnecessary rebuilds on subsequent cmake runs.
file(WRITE ${dir}/dummy.cpp
"#include <Test${Tag}_Category.hpp>\n"
"#include <${SOURCES_A_Input}.hpp>\n"
"#include <${SOURCE_Input}.hpp>\n"
)
configure_file(${dir}/dummy.cpp ${file})
list(APPEND SOURCES_A ${file})
list(APPEND ALGO_SORT_SOURCES ${file})
endforeach()
# ------------------------------------------
# Random
# ------------------------------------------
# do as above
set(ALGO_RANDOM_SOURCES)
foreach(SOURCE_Input
TestRandom
)
set(file ${dir}/${SOURCE_Input}.cpp)
file(WRITE ${dir}/dummy.cpp
"#include <Test${Tag}_Category.hpp>\n"
"#include <${SOURCE_Input}.hpp>\n"
)
configure_file(${dir}/dummy.cpp ${file})
list(APPEND ALGO_RANDOM_SOURCES ${file})
endforeach()
# ------------------------------------------
@ -145,6 +155,26 @@ foreach(Tag Threads;Serial;OpenMP;Cuda;HPX;HIP;SYCL;OpenMPTarget)
endif()
endforeach()
# FIXME_OPENMPTARGET This test causes internal compiler errors as of 09/01/22
# when compiling for Intel's Xe-HP GPUs.
# FRIZZI: 04/26/2023: not sure if the compilation error is still applicable
# but we conservatively leave this guard on
if(NOT (KOKKOS_ENABLE_OPENMPTARGET AND KOKKOS_CXX_COMPILER_ID STREQUAL IntelLLVM))
KOKKOS_ADD_EXECUTABLE_AND_TEST(
UnitTest_Sort
SOURCES
UnitTestMain.cpp
${ALGO_SORT_SOURCES}
)
KOKKOS_ADD_EXECUTABLE_AND_TEST(
UnitTest_Random
SOURCES
UnitTestMain.cpp
${ALGO_RANDOM_SOURCES}
)
endif()
# FIXME_OPENMPTARGET These tests cause internal compiler errors as of 09/01/22
# when compiling for Intel's Xe-HP GPUs.
if(KOKKOS_ENABLE_OPENMPTARGET AND KOKKOS_CXX_COMPILER_ID STREQUAL IntelLLVM)
@ -160,20 +190,9 @@ if(KOKKOS_ENABLE_OPENMPTARGET AND KOKKOS_CXX_COMPILER_ID STREQUAL IntelLLVM)
)
endif()
# FIXME_OPENMPTARGET This test causes internal compiler errors as of 09/01/22
# when compiling for Intel's Xe-HP GPUs.
if(NOT (KOKKOS_ENABLE_OPENMPTARGET AND KOKKOS_CXX_COMPILER_ID STREQUAL IntelLLVM))
KOKKOS_ADD_EXECUTABLE_AND_TEST(
UnitTest_RandomAndSort
SOURCES
UnitTestMain.cpp
${SOURCES_A}
)
endif()
foreach(ID A;B;C;D;E)
KOKKOS_ADD_EXECUTABLE_AND_TEST(
UnitTest_StdSet_${ID}
AlgorithmsUnitTest_StdSet_${ID}
SOURCES
UnitTestMain.cpp
${STDALGO_SOURCES_${ID}}
@ -184,7 +203,7 @@ endforeach()
# when compiling for Intel's Xe-HP GPUs.
if(NOT (KOKKOS_ENABLE_OPENMPTARGET AND KOKKOS_CXX_COMPILER_ID STREQUAL IntelLLVM))
KOKKOS_ADD_EXECUTABLE(
UnitTest_StdAlgoCompileOnly
AlgorithmsUnitTest_StdAlgoCompileOnly
SOURCES TestStdAlgorithmsCompileOnly.cpp
)
endif()

8
lib/kokkos/algorithms/unit_tests/Makefile
View File

@ -27,10 +27,8 @@ TARGETS =
tmp := $(foreach device, $(KOKKOS_DEVICELIST), \
$(if $(filter Test$(device).cpp, $(shell ls Test$(device).cpp 2>/dev/null)),,\
$(shell echo "\#include <Test"${device}"_Category.hpp>" > Test$(device).cpp); \
$(shell echo "\#include <TestRandomCommon.hpp>" >> Test$(device).cpp); \
$(shell echo "\#include <TestSortCommon.hpp>" >> Test$(device).cpp); \
) \
$(shell echo "\#include <Test"${device}"_Category.hpp>" > Test$(device).cpp); \
) \
)
ifeq ($(KOKKOS_INTERNAL_USE_CUDA), 1)
@ -52,7 +50,7 @@ ifeq ($(KOKKOS_INTERNAL_USE_THREADS), 1)
endif
ifeq ($(KOKKOS_INTERNAL_USE_OPENMP), 1)
OBJ_OPENMP = TestOpenMP.o TestOpenMP_Sort1D.o TestOpenMP_Sort3D.o TestOpenMP_SortDynamicView.o UnitTestMain.o gtest-all.o
OBJ_OPENMP = TestOpenMP.o UnitTestMain.o gtest-all.o
TARGETS += KokkosAlgorithms_UnitTest_OpenMP
TEST_TARGETS += test-openmp
endif

280
lib/kokkos/algorithms/unit_tests/TestBinSortA.hpp Normal file
View File

@ -0,0 +1,280 @@
//@HEADER
// ************************************************************************
//
// Kokkos v. 4.0
// Copyright (2022) National Technology & Engineering
// Solutions of Sandia, LLC (NTESS).
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// the U.S. Government retains certain rights in this software.
//
// Part of Kokkos, under the Apache License v2.0 with LLVM Exceptions.
// See https://kokkos.org/LICENSE for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//@HEADER
#ifndef KOKKOS_ALGORITHMS_UNITTESTS_TEST_BINSORTA_HPP
#define KOKKOS_ALGORITHMS_UNITTESTS_TEST_BINSORTA_HPP
#include <gtest/gtest.h>
#include <Kokkos_Core.hpp>
#include <Kokkos_Random.hpp>
#include <Kokkos_Sort.hpp>
#include <random>
namespace Test {
namespace BinSortSetA {
template <class ExecutionSpace, class Scalar>
struct bin3d_is_sorted_struct {
using value_type = unsigned int;
using execution_space = ExecutionSpace;
Kokkos::View<Scalar * [3], ExecutionSpace> keys;
int max_bins;
Scalar min;
Scalar max;
bin3d_is_sorted_struct(Kokkos::View<Scalar * [3], ExecutionSpace> keys_,
int max_bins_, Scalar min_, Scalar max_)
: keys(keys_), max_bins(max_bins_), min(min_), max(max_) {}
KOKKOS_INLINE_FUNCTION
void operator()(int i, unsigned int& count) const {
int ix1 = int((keys(i, 0) - min) / max * max_bins);
int iy1 = int((keys(i, 1) - min) / max * max_bins);
int iz1 = int((keys(i, 2) - min) / max * max_bins);
int ix2 = int((keys(i + 1, 0) - min) / max * max_bins);
int iy2 = int((keys(i + 1, 1) - min) / max * max_bins);
int iz2 = int((keys(i + 1, 2) - min) / max * max_bins);
if (ix1 > ix2)
count++;
else if (ix1 == ix2) {
if (iy1 > iy2)
count++;
else if ((iy1 == iy2) && (iz1 > iz2))
count++;
}
}
};
template <class ExecutionSpace, class Scalar>
struct sum3D {
using value_type = double;
using execution_space = ExecutionSpace;
Kokkos::View<Scalar * [3], ExecutionSpace> keys;
sum3D(Kokkos::View<Scalar * [3], ExecutionSpace> keys_) : keys(keys_) {}
KOKKOS_INLINE_FUNCTION
void operator()(int i, double& count) const {
count += keys(i, 0);
count += keys(i, 1);
count += keys(i, 2);
}
};
template <class ExecutionSpace, typename KeyType>
void test_3D_sort_impl(unsigned int n) {
using KeyViewType = Kokkos::View<KeyType * [3], ExecutionSpace>;
KeyViewType keys("Keys", n * n * n);
Kokkos::Random_XorShift64_Pool<ExecutionSpace> g(1931);
Kokkos::fill_random(keys, g, 100.0);
double sum_before = 0.0;
double sum_after = 0.0;
unsigned int sort_fails = 0;
ExecutionSpace exec;
Kokkos::parallel_reduce(
Kokkos::RangePolicy<ExecutionSpace>(exec, 0, keys.extent(0)),
sum3D<ExecutionSpace, KeyType>(keys), sum_before);
int bin_1d = 1;
while (bin_1d * bin_1d * bin_1d * 4 < (int)keys.extent(0)) bin_1d *= 2;
int bin_max[3] = {bin_1d, bin_1d, bin_1d};
typename KeyViewType::value_type min[3] = {0, 0, 0};
typename KeyViewType::value_type max[3] = {100, 100, 100};
using BinOp = Kokkos::BinOp3D<KeyViewType>;
BinOp bin_op(bin_max, min, max);
Kokkos::BinSort<KeyViewType, BinOp> Sorter(keys, bin_op, false);
Sorter.create_permute_vector(exec);
Sorter.sort(exec, keys);
Kokkos::parallel_reduce(
Kokkos::RangePolicy<ExecutionSpace>(exec, 0, keys.extent(0)),
sum3D<ExecutionSpace, KeyType>(keys), sum_after);
Kokkos::parallel_reduce(
Kokkos::RangePolicy<ExecutionSpace>(exec, 0, keys.extent(0) - 1),
bin3d_is_sorted_struct<ExecutionSpace, KeyType>(keys, bin_1d, min[0],
max[0]),
sort_fails);
double ratio = sum_before / sum_after;
double epsilon = 1e-10;
unsigned int equal_sum =
(ratio > (1.0 - epsilon)) && (ratio < (1.0 + epsilon)) ? 1 : 0;
if (sort_fails)
printf("3D Sort Sum: %f %f Fails: %u\n", sum_before, sum_after, sort_fails);
ASSERT_EQ(sort_fails, 0u);
ASSERT_EQ(equal_sum, 1u);
}
template <class ExecutionSpace>
void test_issue_1160_impl() {
Kokkos::View<int*, ExecutionSpace> element_("element", 10);
Kokkos::View<double*, ExecutionSpace> x_("x", 10);
Kokkos::View<double*, ExecutionSpace> v_("y", 10);
auto h_element = Kokkos::create_mirror_view(element_);
auto h_x = Kokkos::create_mirror_view(x_);
auto h_v = Kokkos::create_mirror_view(v_);
h_element(0) = 9;
h_element(1) = 8;
h_element(2) = 7;
h_element(3) = 6;
h_element(4) = 5;
h_element(5) = 4;
h_element(6) = 3;
h_element(7) = 2;
h_element(8) = 1;
h_element(9) = 0;
for (int i = 0; i < 10; ++i) {
h_v.access(i, 0) = h_x.access(i, 0) = double(h_element(i));
}
ExecutionSpace exec;
Kokkos::deep_copy(exec, element_, h_element);
Kokkos::deep_copy(exec, x_, h_x);
Kokkos::deep_copy(exec, v_, h_v);
using KeyViewType = decltype(element_);
using BinOp = Kokkos::BinOp1D<KeyViewType>;
int begin = 3;
int end = 8;
auto max = h_element(begin);
auto min = h_element(end - 1);
BinOp binner(end - begin, min, max);
Kokkos::BinSort<KeyViewType, BinOp> Sorter(element_, begin, end, binner,
false);
Sorter.create_permute_vector(exec);
Sorter.sort(exec, element_, begin, end);
Sorter.sort(exec, x_, begin, end);
Sorter.sort(exec, v_, begin, end);
Kokkos::deep_copy(exec, h_element, element_);
Kokkos::deep_copy(exec, h_x, x_);
Kokkos::deep_copy(exec, h_v, v_);
exec.fence();
ASSERT_EQ(h_element(0), 9);
ASSERT_EQ(h_element(1), 8);
ASSERT_EQ(h_element(2), 7);
ASSERT_EQ(h_element(3), 2);
ASSERT_EQ(h_element(4), 3);
ASSERT_EQ(h_element(5), 4);
ASSERT_EQ(h_element(6), 5);
ASSERT_EQ(h_element(7), 6);
ASSERT_EQ(h_element(8), 1);
ASSERT_EQ(h_element(9), 0);
for (int i = 0; i < 10; ++i) {
ASSERT_EQ(h_element(i), int(h_x.access(i, 0)));
ASSERT_EQ(h_element(i), int(h_v.access(i, 0)));
}
}
template <class ExecutionSpace, class T>
void test_sort_integer_overflow() {
// FIXME: this test is meant to test something for BinSort,
// but actually uses the kokkos::sort API with the assumption
// that underneath it calls binsort. I don't think this is correct,
// because if the kokkos::sort API chages impl, this test is not testing
// what it meants to test... so need to change this to actually use BinSort
// directly.
// array with two extrema in reverse order to expose integer overflow bug in
// bin calculation
T a[2] = {Kokkos::Experimental::finite_max<T>::value,
Kokkos::Experimental::finite_min<T>::value};
auto vd = Kokkos::create_mirror_view_and_copy(
ExecutionSpace(), Kokkos::View<T[2], Kokkos::HostSpace>(a));
Kokkos::sort(vd);
auto vh = Kokkos::create_mirror_view_and_copy(Kokkos::HostSpace(), vd);
EXPECT_TRUE(std::is_sorted(vh.data(), vh.data() + 2))
<< "view (" << vh[0] << ", " << vh[1] << ") is not sorted";
}
} // namespace BinSortSetA
TEST(TEST_CATEGORY, BinSortGenericTests) {
using ExecutionSpace = TEST_EXECSPACE;
using key_type = unsigned;
constexpr int N = 171;
#if defined(KOKKOS_ENABLE_CUDA) && \
defined(KOKKOS_COMPILER_NVHPC) // FIXME_NVHPC
if (!std::is_same_v<ExecutionSpace, Kokkos::Cuda>)
#endif
BinSortSetA::test_3D_sort_impl<ExecutionSpace, key_type>(N);
#if defined(KOKKOS_ENABLE_CUDA) && \
defined(KOKKOS_COMPILER_NVHPC) // FIXME_NVHPC
if (!std::is_same_v<ExecutionSpace, Kokkos::Cuda>)
#endif
BinSortSetA::test_issue_1160_impl<ExecutionSpace>();
BinSortSetA::test_sort_integer_overflow<ExecutionSpace, long long>();
BinSortSetA::test_sort_integer_overflow<ExecutionSpace, unsigned long long>();
BinSortSetA::test_sort_integer_overflow<ExecutionSpace, int>();
}
TEST(TEST_CATEGORY, BinSortEmptyView) {
using ExecutionSpace = TEST_EXECSPACE;
// the bounds and extents used below are totally arbitrary
// and, in theory, should have no impact
using KeyViewType = Kokkos::View<int*, ExecutionSpace>;
KeyViewType kv("kv", 20);
using BinOp_t = Kokkos::BinOp1D<KeyViewType>;
BinOp_t binOp(5, 0, 10);
Kokkos::BinSort<KeyViewType, BinOp_t> Sorter(ExecutionSpace{}, kv, binOp);
// does not matter if we use int or something else
Kokkos::View<int*, ExecutionSpace> v("v", 0);
// test all exposed public sort methods
ASSERT_NO_THROW(Sorter.sort(ExecutionSpace(), v, 0, 0));
ASSERT_NO_THROW(Sorter.sort(v, 0, 0));
ASSERT_NO_THROW(Sorter.sort(ExecutionSpace(), v));
ASSERT_NO_THROW(Sorter.sort(v));
}
TEST(TEST_CATEGORY, BinSortEmptyKeysView) {
using ExecutionSpace = TEST_EXECSPACE;
using KeyViewType = Kokkos::View<int*, ExecutionSpace>;
KeyViewType kv("kv", 0);
using BinOp_t = Kokkos::BinOp1D<KeyViewType>;
BinOp_t binOp(5, 0, 10);
Kokkos::BinSort<KeyViewType, BinOp_t> Sorter(ExecutionSpace{}, kv, binOp);
ASSERT_NO_THROW(Sorter.create_permute_vector(ExecutionSpace{}));
}
} // namespace Test
#endif

262
lib/kokkos/algorithms/unit_tests/TestBinSortB.hpp Normal file
View File

@ -0,0 +1,262 @@
//@HEADER
// ************************************************************************
//
// Kokkos v. 4.0
// Copyright (2022) National Technology & Engineering
// Solutions of Sandia, LLC (NTESS).
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// the U.S. Government retains certain rights in this software.
//
// Part of Kokkos, under the Apache License v2.0 with LLVM Exceptions.
// See https://kokkos.org/LICENSE for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//@HEADER
#ifndef KOKKOS_ALGORITHMS_UNITTESTS_TEST_BINSORTB_HPP
#define KOKKOS_ALGORITHMS_UNITTESTS_TEST_BINSORTB_HPP
#include <gtest/gtest.h>
#include <Kokkos_Core.hpp>
#include <Kokkos_Random.hpp>
#include <Kokkos_Sort.hpp>
#include <Kokkos_StdAlgorithms.hpp>
#include <TestStdAlgorithmsHelperFunctors.hpp>
#include <random>
#include <numeric> //needed for iota
namespace Test {
namespace BinSortSetB {
template <class ViewTypeFrom, class ViewTypeTo>
struct CopyFunctorRank2 {
ViewTypeFrom m_view_from;
ViewTypeTo m_view_to;
CopyFunctorRank2() = delete;
CopyFunctorRank2(const ViewTypeFrom view_from, const ViewTypeTo view_to)
: m_view_from(view_from), m_view_to(view_to) {}
KOKKOS_INLINE_FUNCTION
void operator()(int k) const {
const auto i = k / m_view_from.extent(1);
const auto j = k % m_view_from.extent(1);
m_view_to(i, j) = m_view_from(i, j);
}
};
template <class ValueType, class... Props>
auto create_deep_copyable_compatible_view_with_same_extent(
Kokkos::View<ValueType*, Props...> view) {
using view_type = Kokkos::View<ValueType*, Props...>;
using view_value_type = typename view_type::value_type;
using view_exespace = typename view_type::execution_space;
const std::size_t ext0 = view.extent(0);
using view_deep_copyable_t = Kokkos::View<view_value_type*, view_exespace>;
return view_deep_copyable_t{"view_dc", ext0};
}
template <class ValueType, class... Props>
auto create_deep_copyable_compatible_view_with_same_extent(
Kokkos::View<ValueType**, Props...> view) {
using view_type = Kokkos::View<ValueType**, Props...>;
using view_value_type = typename view_type::value_type;
using view_exespace = typename view_type::execution_space;
using view_deep_copyable_t = Kokkos::View<view_value_type**, view_exespace>;
const std::size_t ext0 = view.extent(0);
const std::size_t ext1 = view.extent(1);
return view_deep_copyable_t{"view_dc", ext0, ext1};
}
template <class ViewType>
auto create_deep_copyable_compatible_clone(ViewType view) {
static_assert(ViewType::rank <= 2);
auto view_dc = create_deep_copyable_compatible_view_with_same_extent(view);
using view_dc_t = decltype(view_dc);
if constexpr (ViewType::rank == 1) {
Test::stdalgos::CopyFunctor<ViewType, view_dc_t> F1(view, view_dc);
Kokkos::parallel_for("copy", view.extent(0), F1);
} else {
static_assert(ViewType::rank == 2, "Only rank 1 or 2 supported.");
CopyFunctorRank2<ViewType, view_dc_t> F1(view, view_dc);
Kokkos::parallel_for("copy", view.extent(0) * view.extent(1), F1);
}
return view_dc;
}
template <class ViewType>
auto create_host_space_copy(ViewType view) {
auto view_dc = create_deep_copyable_compatible_clone(view);
return create_mirror_view_and_copy(Kokkos::HostSpace(), view_dc);
}
template <class KeyType, class ExecutionSpace>
auto create_rank1_dev_and_host_views_of_keys(const ExecutionSpace& exec,
int N) {
namespace KE = Kokkos::Experimental;
Kokkos::DefaultHostExecutionSpace defaultHostExeSpace;
using KeyViewType = Kokkos::View<KeyType*, ExecutionSpace>;
KeyViewType keys("keys", N);
auto keys_h = Kokkos::create_mirror_view(keys);
std::iota(KE::begin(keys_h), KE::end(keys_h), KeyType(0));
KE::reverse(defaultHostExeSpace, keys_h);
// keys now is = [N-1,N-2,...,2,1,0], shuffle it for avoid trivial case
std::random_device rd;
std::mt19937 g(rd());
std::shuffle(KE::begin(keys_h), KE::end(keys_h), g);
Kokkos::deep_copy(exec, keys, keys_h);
return std::make_pair(keys, keys_h);
}
template <class ExecutionSpace, class ValueType, int ValuesViewRank,
std::enable_if_t<ValuesViewRank == 1, int> = 0>
auto create_strided_view(std::size_t numRows, std::size_t /*numCols*/) {
Kokkos::LayoutStride layout{numRows, 2};
using v_t = Kokkos::View<ValueType*, Kokkos::LayoutStride, ExecutionSpace>;
v_t v("v", layout);
return v;
}
template <class ExecutionSpace, class ValueType, int ValuesViewRank,
std::enable_if_t<ValuesViewRank == 2, int> = 0>
auto create_strided_view(std::size_t numRows, std::size_t numCols) {
Kokkos::LayoutStride layout{numRows, 2, numCols, numRows * 2};
using v_t = Kokkos::View<ValueType**, Kokkos::LayoutStride, ExecutionSpace>;
v_t v("v", layout);
return v;
}
template <class ExecutionSpace, class KeyType, class ValueType,
int ValuesViewRank>
void test_on_view_with_stride(std::size_t numRows, std::size_t indB,
std::size_t indE, std::size_t numCols = 1) {
ExecutionSpace exec;
Kokkos::DefaultHostExecutionSpace defaultHostExeSpace;
namespace KE = Kokkos::Experimental;
// 1. generate 1D view of keys
auto [keys, keys_h] =
create_rank1_dev_and_host_views_of_keys<KeyType>(exec, numRows);
using KeyViewType = decltype(keys);
// need this map key->row to use later for checking
std::unordered_map<KeyType, std::size_t> keyToRowBeforeSort;
for (std::size_t i = 0; i < numRows; ++i) {
keyToRowBeforeSort[keys_h(i)] = i;
}
// 2. create binOp
using BinOp = Kokkos::BinOp1D<KeyViewType>;
auto itB = KE::cbegin(keys_h) + indB;
auto itE = itB + indE - indB;
auto it = KE::minmax_element(defaultHostExeSpace, itB, itE);
// seems like the behavior is odd when we use # buckets = # keys
// so use +5 for using more buckets than keys.
// This is something to investigate.
BinOp binner(indE - indB + 5, *it.first, *it.second);
// 3. create sorter
Kokkos::BinSort<KeyViewType, BinOp> sorter(keys, indB, indE, binner, false);
sorter.create_permute_vector(exec);
sorter.sort(exec, keys, indB, indE);
Kokkos::deep_copy(exec, keys_h, keys);
auto v = create_strided_view<ExecutionSpace, ValueType, ValuesViewRank>(
numRows, numCols);
Kokkos::Random_XorShift64_Pool<ExecutionSpace> pool(73931);
Kokkos::fill_random(v, pool, ValueType(545));
auto v_before_sort_h = create_host_space_copy(v);
sorter.sort(exec, v, indB, indE);
auto v_after_sort_h = create_host_space_copy(v);
for (size_t i = 0; i < v.extent(0); ++i) {
// if i within [indB,indE), the sorting was done
// so we need to do proper checking since rows have changed
if (i >= size_t(indB) && i < size_t(indE)) {
const KeyType key = keys_h(i);
if constexpr (ValuesViewRank == 1) {
ASSERT_TRUE(v_before_sort_h(keyToRowBeforeSort.at(key)) ==
v_after_sort_h(i));
} else {
for (size_t j = 0; j < v.extent(1); ++j) {
ASSERT_TRUE(v_before_sort_h(keyToRowBeforeSort.at(key), j) ==
v_after_sort_h(i, j));
}
}
}
// outside the target bounds, then the i-th row remains unchanged
else {
if constexpr (ValuesViewRank == 1) {
ASSERT_TRUE(v_before_sort_h(i) == v_after_sort_h(i));
} else {
for (size_t j = 0; j < v.extent(1); ++j) {
ASSERT_TRUE(v_before_sort_h(i, j) == v_after_sort_h(i, j));
}
}
}
}
}
template <class ExecutionSpace, class KeyType, class ValueType>
void run_for_rank1() {
constexpr int rank = 1;
// trivial case
test_on_view_with_stride<ExecutionSpace, KeyType, ValueType, rank>(1, 0, 1);
// nontrivial cases
for (std::size_t N : {311, 710017}) {
// various cases for bounds
test_on_view_with_stride<ExecutionSpace, KeyType, ValueType, rank>(N, 0, N);
test_on_view_with_stride<ExecutionSpace, KeyType, ValueType, rank>(N, 3, N);
test_on_view_with_stride<ExecutionSpace, KeyType, ValueType, rank>(N, 0,
N - 4);
test_on_view_with_stride<ExecutionSpace, KeyType, ValueType, rank>(N, 4,
N - 3);
}
}
template <class ExecutionSpace, class KeyType, class ValueType>
void run_for_rank2() {
constexpr int rank = 2;
// trivial case
test_on_view_with_stride<ExecutionSpace, KeyType, ValueType, rank>(1, 0, 1,
1);
// nontrivial cases
for (std::size_t Nr : {11, 1157, 710017}) {
for (std::size_t Nc : {3, 51}) {
// various cases for bounds
test_on_view_with_stride<ExecutionSpace, KeyType, ValueType, rank>(
Nr, 0, Nr, Nc);
test_on_view_with_stride<ExecutionSpace, KeyType, ValueType, rank>(
Nr, 3, Nr, Nc);
test_on_view_with_stride<ExecutionSpace, KeyType, ValueType, rank>(
Nr, 0, Nr - 4, Nc);
test_on_view_with_stride<ExecutionSpace, KeyType, ValueType, rank>(
Nr, 4, Nr - 3, Nc);
}
}
}
} // namespace BinSortSetB
TEST(TEST_CATEGORY, BinSortUnsignedKeyLayoutStrideValues) {
using ExeSpace = TEST_EXECSPACE;
using key_type = unsigned;
BinSortSetB::run_for_rank1<ExeSpace, key_type, int>();
BinSortSetB::run_for_rank1<ExeSpace, key_type, double>();
BinSortSetB::run_for_rank2<ExeSpace, key_type, int>();
BinSortSetB::run_for_rank2<ExeSpace, key_type, double>();
}
} // namespace Test
#endif

22
lib/kokkos/algorithms/unit_tests/TestNestedSort.hpp
View File

@ -17,14 +17,14 @@
#ifndef KOKKOS_ALGORITHMS_UNITTESTS_TEST_NESTED_SORT_HPP
#define KOKKOS_ALGORITHMS_UNITTESTS_TEST_NESTED_SORT_HPP
#include <gtest/gtest.h>
#include <unordered_set>
#include <random>
#include <Kokkos_Random.hpp>
#include <Kokkos_NestedSort.hpp>
namespace Test {
namespace Impl {
namespace NestedSortImpl {
// Comparator for sorting in descending order
template <typename Key>
@ -383,24 +383,28 @@ void test_nested_sort_by_key(unsigned int N, KeyType minKey, KeyType maxKey,
test_nested_sort_by_key_impl<ExecutionSpace, KeyType, ValueType>(
N, N, false, true, minKey, maxKey, minVal, maxVal);
}
} // namespace Impl
} // namespace NestedSortImpl
TEST(TEST_CATEGORY, NestedSort) {
Impl::test_nested_sort<TEST_EXECSPACE, unsigned>(171, 0U, UINT_MAX);
Impl::test_nested_sort<TEST_EXECSPACE, float>(42, -1e6f, 1e6f);
Impl::test_nested_sort<TEST_EXECSPACE, char>(67, CHAR_MIN, CHAR_MAX);
using ExecutionSpace = TEST_EXECSPACE;
NestedSortImpl::test_nested_sort<ExecutionSpace, unsigned>(171, 0U, UINT_MAX);
NestedSortImpl::test_nested_sort<ExecutionSpace, float>(42, -1e6f, 1e6f);
NestedSortImpl::test_nested_sort<ExecutionSpace, char>(67, CHAR_MIN,
CHAR_MAX);
}
TEST(TEST_CATEGORY, NestedSortByKey) {
using ExecutionSpace = TEST_EXECSPACE;
// Second/third template arguments are key and value respectively.
// In sort_by_key_X functions, a key view and a value view are both permuted
// to make the keys sorted. This means that the value type doesn't need to be
// ordered, unlike key
Impl::test_nested_sort_by_key<TEST_EXECSPACE, unsigned, unsigned>(
NestedSortImpl::test_nested_sort_by_key<ExecutionSpace, unsigned, unsigned>(
161, 0U, UINT_MAX, 0U, UINT_MAX);
Impl::test_nested_sort_by_key<TEST_EXECSPACE, float, char>(
NestedSortImpl::test_nested_sort_by_key<ExecutionSpace, float, char>(
267, -1e6f, 1e6f, CHAR_MIN, CHAR_MAX);
Impl::test_nested_sort_by_key<TEST_EXECSPACE, char, double>(
NestedSortImpl::test_nested_sort_by_key<ExecutionSpace, char, double>(
11, CHAR_MIN, CHAR_MAX, 2.718, 3.14);
}

39
lib/kokkos/algorithms/unit_tests/TestOpenMP_Sort1D.cpp
View File

@ -1,39 +0,0 @@
//@HEADER
// ************************************************************************
//
// Kokkos v. 4.0
// Copyright (2022) National Technology & Engineering
// Solutions of Sandia, LLC (NTESS).
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// the U.S. Government retains certain rights in this software.
//
// Part of Kokkos, under the Apache License v2.0 with LLVM Exceptions.
// See https://kokkos.org/LICENSE for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//@HEADER
#include <Kokkos_Macros.hpp>
#ifdef KOKKOS_ENABLE_OPENMP
#include <gtest/gtest.h>
#include <Kokkos_Core.hpp>
//----------------------------------------------------------------------------
#include <TestRandom.hpp>
#include <TestSort.hpp>
#include <iomanip>
namespace Test {
TEST(openmp, SortUnsigned1D) {
Impl::test_1D_sort<Kokkos::OpenMP, unsigned>(171);
}
TEST(openmp, SortIssue1160) { Impl::test_issue_1160_sort<Kokkos::OpenMP>(); }
} // namespace Test
#else
void KOKKOS_ALGORITHMS_UNITTESTS_TESTOPENMP_PREVENT_LINK_ERROR() {}
#endif

37
lib/kokkos/algorithms/unit_tests/TestOpenMP_Sort3D.cpp
View File

@ -1,37 +0,0 @@
//@HEADER
// ************************************************************************
//
// Kokkos v. 4.0
// Copyright (2022) National Technology & Engineering
// Solutions of Sandia, LLC (NTESS).
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// the U.S. Government retains certain rights in this software.
//
// Part of Kokkos, under the Apache License v2.0 with LLVM Exceptions.
// See https://kokkos.org/LICENSE for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//@HEADER
#include <Kokkos_Macros.hpp>
#ifdef KOKKOS_ENABLE_OPENMP
#include <gtest/gtest.h>
#include <Kokkos_Core.hpp>
//----------------------------------------------------------------------------
#include <TestRandom.hpp>
#include <TestSort.hpp>
#include <iomanip>
namespace Test {
TEST(openmp, SortUnsigned3D) {
Impl::test_3D_sort<Kokkos::OpenMP, unsigned>(171);
}
} // namespace Test
#else
void KOKKOS_ALGORITHMS_UNITTESTS_TESTOPENMP_PREVENT_LINK_ERROR() {}
#endif

37
lib/kokkos/algorithms/unit_tests/TestOpenMP_SortDynamicView.cpp
View File

@ -1,37 +0,0 @@
//@HEADER
// ************************************************************************
//
// Kokkos v. 4.0
// Copyright (2022) National Technology & Engineering
// Solutions of Sandia, LLC (NTESS).
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// the U.S. Government retains certain rights in this software.
//
// Part of Kokkos, under the Apache License v2.0 with LLVM Exceptions.
// See https://kokkos.org/LICENSE for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//@HEADER
#include <Kokkos_Macros.hpp>
#ifdef KOKKOS_ENABLE_OPENMP
#include <gtest/gtest.h>
#include <Kokkos_Core.hpp>
//----------------------------------------------------------------------------
#include <TestRandom.hpp>
#include <TestSort.hpp>
#include <iomanip>
namespace Test {
TEST(openmp, SortUnsignedDynamicView) {
Impl::test_dynamic_view_sort<Kokkos::OpenMP, unsigned>(171);
}
} // namespace Test
#else
void KOKKOS_ALGORITHMS_UNITTESTS_TESTOPENMP_PREVENT_LINK_ERROR() {}
#endif

41
lib/kokkos/algorithms/unit_tests/TestRandom.hpp
View File

@ -14,8 +14,8 @@
//
//@HEADER
#ifndef KOKKOS_TEST_DUALVIEW_HPP
#define KOKKOS_TEST_DUALVIEW_HPP
#ifndef KOKKOS_ALGORITHMS_UNITTESTS_TEST_RANDOM_HPP
#define KOKKOS_ALGORITHMS_UNITTESTS_TEST_RANDOM_HPP
#include <gtest/gtest.h>
#include <iostream>
@ -29,8 +29,7 @@
#include <chrono>
namespace Test {
namespace Impl {
namespace AlgoRandomImpl {
// This test runs the random number generators and uses some statistic tests to
// check the 'goodness' of the random numbers:
@ -469,42 +468,46 @@ struct TestDynRankView {
ASSERT_LE(val.max_val, max);
}
};
} // namespace Impl
template <typename ExecutionSpace>
void test_random_xorshift64() {
} // namespace AlgoRandomImpl
TEST(TEST_CATEGORY, Random_XorShift64) {
using ExecutionSpace = TEST_EXECSPACE;
#if defined(KOKKOS_ENABLE_SYCL) || defined(KOKKOS_ENABLE_CUDA) || \
defined(KOKKOS_ENABLE_HIP)
const int num_draws = 132141141;
#else // SERIAL, HPX, OPENMP
const int num_draws = 10240000;
#endif
Impl::test_random<Kokkos::Random_XorShift64_Pool<ExecutionSpace>>(num_draws);
Impl::test_random<Kokkos::Random_XorShift64_Pool<
AlgoRandomImpl::test_random<Kokkos::Random_XorShift64_Pool<ExecutionSpace>>(
num_draws);
AlgoRandomImpl::test_random<Kokkos::Random_XorShift64_Pool<
Kokkos::Device<ExecutionSpace, typename ExecutionSpace::memory_space>>>(
num_draws);
Impl::TestDynRankView<ExecutionSpace,
Kokkos::Random_XorShift64_Pool<ExecutionSpace>>(10000)
AlgoRandomImpl::TestDynRankView<
ExecutionSpace, Kokkos::Random_XorShift64_Pool<ExecutionSpace>>(10000)
.run();
}
template <typename ExecutionSpace>
void test_random_xorshift1024() {
TEST(TEST_CATEGORY, Random_XorShift1024_0) {
using ExecutionSpace = TEST_EXECSPACE;
#if defined(KOKKOS_ENABLE_SYCL) || defined(KOKKOS_ENABLE_CUDA) || \
defined(KOKKOS_ENABLE_HIP)
const int num_draws = 52428813;
#else // SERIAL, HPX, OPENMP
const int num_draws = 10130144;
#endif
Impl::test_random<Kokkos::Random_XorShift1024_Pool<ExecutionSpace>>(
AlgoRandomImpl::test_random<Kokkos::Random_XorShift1024_Pool<ExecutionSpace>>(
num_draws);
Impl::test_random<Kokkos::Random_XorShift1024_Pool<
AlgoRandomImpl::test_random<Kokkos::Random_XorShift1024_Pool<
Kokkos::Device<ExecutionSpace, typename ExecutionSpace::memory_space>>>(
num_draws);
Impl::TestDynRankView<ExecutionSpace,
Kokkos::Random_XorShift1024_Pool<ExecutionSpace>>(10000)
AlgoRandomImpl::TestDynRankView<
ExecutionSpace, Kokkos::Random_XorShift1024_Pool<ExecutionSpace>>(10000)
.run();
}
} // namespace Test
#endif // KOKKOS_TEST_UNORDERED_MAP_HPP
} // namespace Test
#endif

22
lib/kokkos/algorithms/unit_tests/TestRandomAccessIterator.cpp
View File

@ -54,7 +54,7 @@ void test_random_access_it_verify(IteratorType it, ValueType gold_value) {
Kokkos::parallel_for("_std_algo_copy", 1, cf);
auto v_h =
Kokkos::create_mirror_view_and_copy(Kokkos::HostSpace(), checkView);
EXPECT_EQ(v_h(), gold_value);
ASSERT_EQ(v_h(), gold_value);
}
TEST_F(random_access_iterator_test, dereference) {
@ -96,9 +96,9 @@ void test_random_access_it_subscript_op_verify(IteratorType it) {
auto v_h =
Kokkos::create_mirror_view_and_copy(Kokkos::HostSpace(), checkView);
EXPECT_EQ(v_h(0), (value_t)0);
EXPECT_EQ(v_h(1), (value_t)1);
EXPECT_EQ(v_h(2), (value_t)2);
ASSERT_EQ(v_h(0), (value_t)0);
ASSERT_EQ(v_h(1), (value_t)1);
ASSERT_EQ(v_h(2), (value_t)2);
}
TEST_F(random_access_iterator_test, subscript_operator) {
@ -188,9 +188,9 @@ TEST_F(random_access_iterator_test, operatorsSet4) {
auto it7 = KE::Impl::RandomAccessIterator<static_view_t>(m_static_view, 3);
auto it8 = KE::Impl::RandomAccessIterator<dyn_view_t>(m_dynamic_view, 3);
auto it9 = KE::Impl::RandomAccessIterator<strided_view_t>(m_strided_view, 3);
EXPECT_EQ(it1, it7);
EXPECT_EQ(it2, it8);
EXPECT_EQ(it3, it9);
ASSERT_EQ(it1, it7);
ASSERT_EQ(it2, it8);
ASSERT_EQ(it3, it9);
EXPECT_GE(it1, it7);
EXPECT_GE(it2, it8);
EXPECT_GE(it3, it9);
@ -205,16 +205,16 @@ TEST_F(random_access_iterator_test, assignment_operator) {
EXPECT_NE(it1, it2);
it2 = it1;
EXPECT_EQ(it1, it2);
ASSERT_EQ(it1, it2);
}
TEST_F(random_access_iterator_test, distance) {
auto first = KE::begin(m_dynamic_view);
auto last = KE::end(m_dynamic_view);
EXPECT_EQ(0, KE::distance(first, first));
EXPECT_EQ(1, KE::distance(first, first + 1));
EXPECT_EQ(m_dynamic_view.extent(0), size_t(KE::distance(first, last)));
ASSERT_EQ(0, KE::distance(first, first));
ASSERT_EQ(1, KE::distance(first, first + 1));
ASSERT_EQ(m_dynamic_view.extent(0), size_t(KE::distance(first, last)));
}
} // namespace stdalgos

32
lib/kokkos/algorithms/unit_tests/TestRandomCommon.hpp
View File

@ -1,32 +0,0 @@
//@HEADER
// ************************************************************************
//
// Kokkos v. 4.0
// Copyright (2022) National Technology & Engineering
// Solutions of Sandia, LLC (NTESS).
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// the U.S. Government retains certain rights in this software.
//
// Part of Kokkos, under the Apache License v2.0 with LLVM Exceptions.
// See https://kokkos.org/LICENSE for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//@HEADER
#ifndef KOKKOS_ALGORITHMS_UNITTESTS_TESTRANDOM_COMMON_HPP
#define KOKKOS_ALGORITHMS_UNITTESTS_TESTRANDOM_COMMON_HPP
#include <TestRandom.hpp>
namespace Test {
TEST(TEST_CATEGORY, Random_XorShift64) {
test_random_xorshift64<TEST_EXECSPACE>();
}
TEST(TEST_CATEGORY, Random_XorShift1024_0) {
test_random_xorshift1024<TEST_EXECSPACE>();
}
} // namespace Test
#endif

244
lib/kokkos/algorithms/unit_tests/TestSort.hpp
View File

@ -14,8 +14,8 @@
//
//@HEADER
#ifndef KOKKOS_ALGORITHMS_UNITTESTS_TESTSORT_HPP
#define KOKKOS_ALGORITHMS_UNITTESTS_TESTSORT_HPP
#ifndef KOKKOS_ALGORITHMS_UNITTESTS_TEST_SORT_HPP
#define KOKKOS_ALGORITHMS_UNITTESTS_TEST_SORT_HPP
#include <gtest/gtest.h>
#include <Kokkos_Core.hpp>
@ -24,8 +24,7 @@
#include <Kokkos_Sort.hpp>
namespace Test {
namespace Impl {
namespace SortImpl {
template <class ExecutionSpace, class Scalar>
struct is_sorted_struct {
@ -53,56 +52,6 @@ struct sum {
void operator()(int i, double& count) const { count += keys(i); }
};
template <class ExecutionSpace, class Scalar>
struct bin3d_is_sorted_struct {
using value_type = unsigned int;
using execution_space = ExecutionSpace;
Kokkos::View<Scalar * [3], ExecutionSpace> keys;
int max_bins;
Scalar min;
Scalar max;
bin3d_is_sorted_struct(Kokkos::View<Scalar * [3], ExecutionSpace> keys_,
int max_bins_, Scalar min_, Scalar max_)
: keys(keys_), max_bins(max_bins_), min(min_), max(max_) {}
KOKKOS_INLINE_FUNCTION
void operator()(int i, unsigned int& count) const {
int ix1 = int((keys(i, 0) - min) / max * max_bins);
int iy1 = int((keys(i, 1) - min) / max * max_bins);
int iz1 = int((keys(i, 2) - min) / max * max_bins);
int ix2 = int((keys(i + 1, 0) - min) / max * max_bins);
int iy2 = int((keys(i + 1, 1) - min) / max * max_bins);
int iz2 = int((keys(i + 1, 2) - min) / max * max_bins);
if (ix1 > ix2)
count++;
else if (ix1 == ix2) {
if (iy1 > iy2)
count++;
else if ((iy1 == iy2) && (iz1 > iz2))
count++;
}
}
};
template <class ExecutionSpace, class Scalar>
struct sum3D {
using value_type = double;
using execution_space = ExecutionSpace;
Kokkos::View<Scalar * [3], ExecutionSpace> keys;
sum3D(Kokkos::View<Scalar * [3], ExecutionSpace> keys_) : keys(keys_) {}
KOKKOS_INLINE_FUNCTION
void operator()(int i, double& count) const {
count += keys(i, 0);
count += keys(i, 1);
count += keys(i, 2);
}
};
template <class ExecutionSpace, typename KeyType>
void test_1D_sort_impl(unsigned int n) {
using KeyViewType = Kokkos::View<KeyType*, ExecutionSpace>;
@ -142,57 +91,6 @@ void test_1D_sort_impl(unsigned int n) {
ASSERT_EQ(equal_sum, 1u);
}
template <class ExecutionSpace, typename KeyType>
void test_3D_sort_impl(unsigned int n) {
using KeyViewType = Kokkos::View<KeyType * [3], ExecutionSpace>;
KeyViewType keys("Keys", n * n * n);
Kokkos::Random_XorShift64_Pool<ExecutionSpace> g(1931);
Kokkos::fill_random(keys, g, 100.0);
double sum_before = 0.0;
double sum_after = 0.0;
unsigned int sort_fails = 0;
ExecutionSpace exec;
Kokkos::parallel_reduce(
Kokkos::RangePolicy<ExecutionSpace>(exec, 0, keys.extent(0)),
sum3D<ExecutionSpace, KeyType>(keys), sum_before);
int bin_1d = 1;
while (bin_1d * bin_1d * bin_1d * 4 < (int)keys.extent(0)) bin_1d *= 2;
int bin_max[3] = {bin_1d, bin_1d, bin_1d};
typename KeyViewType::value_type min[3] = {0, 0, 0};
typename KeyViewType::value_type max[3] = {100, 100, 100};
using BinOp = Kokkos::BinOp3D<KeyViewType>;
BinOp bin_op(bin_max, min, max);
Kokkos::BinSort<KeyViewType, BinOp> Sorter(keys, bin_op, false);
Sorter.create_permute_vector(exec);
Sorter.sort(exec, keys);
Kokkos::parallel_reduce(
Kokkos::RangePolicy<ExecutionSpace>(exec, 0, keys.extent(0)),
sum3D<ExecutionSpace, KeyType>(keys), sum_after);
Kokkos::parallel_reduce(
Kokkos::RangePolicy<ExecutionSpace>(exec, 0, keys.extent(0) - 1),
bin3d_is_sorted_struct<ExecutionSpace, KeyType>(keys, bin_1d, min[0],
max[0]),
sort_fails);
double ratio = sum_before / sum_after;
double epsilon = 1e-10;
unsigned int equal_sum =
(ratio > (1.0 - epsilon)) && (ratio < (1.0 + epsilon)) ? 1 : 0;
if (sort_fails)
printf("3D Sort Sum: %f %f Fails: %u\n", sum_before, sum_after, sort_fails);
ASSERT_EQ(sort_fails, 0u);
ASSERT_EQ(equal_sum, 1u);
}
//----------------------------------------------------------------------------
template <class ExecutionSpace, typename KeyType>
@ -259,74 +157,6 @@ void test_dynamic_view_sort_impl(unsigned int n) {
//----------------------------------------------------------------------------
template <class ExecutionSpace>
void test_issue_1160_impl() {
Kokkos::View<int*, ExecutionSpace> element_("element", 10);
Kokkos::View<double*, ExecutionSpace> x_("x", 10);
Kokkos::View<double*, ExecutionSpace> v_("y", 10);
auto h_element = Kokkos::create_mirror_view(element_);
auto h_x = Kokkos::create_mirror_view(x_);
auto h_v = Kokkos::create_mirror_view(v_);
h_element(0) = 9;
h_element(1) = 8;
h_element(2) = 7;
h_element(3) = 6;
h_element(4) = 5;
h_element(5) = 4;
h_element(6) = 3;
h_element(7) = 2;
h_element(8) = 1;
h_element(9) = 0;
for (int i = 0; i < 10; ++i) {
h_v.access(i, 0) = h_x.access(i, 0) = double(h_element(i));
}
ExecutionSpace exec;
Kokkos::deep_copy(exec, element_, h_element);
Kokkos::deep_copy(exec, x_, h_x);
Kokkos::deep_copy(exec, v_, h_v);
using KeyViewType = decltype(element_);
using BinOp = Kokkos::BinOp1D<KeyViewType>;
int begin = 3;
int end = 8;
auto max = h_element(begin);
auto min = h_element(end - 1);
BinOp binner(end - begin, min, max);
Kokkos::BinSort<KeyViewType, BinOp> Sorter(element_, begin, end, binner,
false);
Sorter.create_permute_vector(exec);
Sorter.sort(exec, element_, begin, end);
Sorter.sort(exec, x_, begin, end);
Sorter.sort(exec, v_, begin, end);
Kokkos::deep_copy(exec, h_element, element_);
Kokkos::deep_copy(exec, h_x, x_);
Kokkos::deep_copy(exec, h_v, v_);
exec.fence();
ASSERT_EQ(h_element(0), 9);
ASSERT_EQ(h_element(1), 8);
ASSERT_EQ(h_element(2), 7);
ASSERT_EQ(h_element(3), 2);
ASSERT_EQ(h_element(4), 3);
ASSERT_EQ(h_element(5), 4);
ASSERT_EQ(h_element(6), 5);
ASSERT_EQ(h_element(7), 6);
ASSERT_EQ(h_element(8), 1);
ASSERT_EQ(h_element(9), 0);
for (int i = 0; i < 10; ++i) {
ASSERT_EQ(h_element(i), int(h_x.access(i, 0)));
ASSERT_EQ(h_element(i), int(h_v.access(i, 0)));
}
}
template <class ExecutionSpace>
void test_issue_4978_impl() {
Kokkos::View<long long*, ExecutionSpace> element_("element", 9);
@ -376,55 +206,33 @@ void test_sort_integer_overflow() {
<< "view (" << vh[0] << ", " << vh[1] << ") is not sorted";
}
//----------------------------------------------------------------------------
} // namespace SortImpl
template <class ExecutionSpace, typename KeyType>
void test_1D_sort(unsigned int N) {
test_1D_sort_impl<ExecutionSpace, KeyType>(N * N * N);
}
TEST(TEST_CATEGORY, SortUnsignedValueType) {
using ExecutionSpace = TEST_EXECSPACE;
using key_type = unsigned;
constexpr int N = 171;
template <class ExecutionSpace, typename KeyType>
void test_3D_sort(unsigned int N) {
test_3D_sort_impl<ExecutionSpace, KeyType>(N);
}
SortImpl::test_1D_sort_impl<ExecutionSpace, key_type>(N * N * N);
template <class ExecutionSpace, typename KeyType>
void test_dynamic_view_sort(unsigned int N) {
test_dynamic_view_sort_impl<ExecutionSpace, KeyType>(N * N);
}
template <class ExecutionSpace>
void test_issue_1160_sort() {
test_issue_1160_impl<ExecutionSpace>();
}
template <class ExecutionSpace>
void test_issue_4978_sort() {
test_issue_4978_impl<ExecutionSpace>();
}
template <class ExecutionSpace, typename KeyType>
void test_sort(unsigned int N) {
test_1D_sort<ExecutionSpace, KeyType>(N);
#if defined(KOKKOS_ENABLE_CUDA) && \
defined(KOKKOS_COMPILER_NVHPC) // FIXME_NVHPC
if (!std::is_same_v<ExecutionSpace, Kokkos::Cuda>)
#endif
test_3D_sort<ExecutionSpace, KeyType>(N);
// FIXME_OPENMPTARGET: OpenMPTarget doesn't support DynamicView yet.
#ifndef KOKKOS_ENABLE_OPENMPTARGET
test_dynamic_view_sort<ExecutionSpace, KeyType>(N);
// FIXME_OPENMPTARGET: OpenMPTarget doesn't support DynamicView yet.
SortImpl::test_dynamic_view_sort_impl<ExecutionSpace, key_type>(N * N);
#endif
#if defined(KOKKOS_ENABLE_CUDA) && \
defined(KOKKOS_COMPILER_NVHPC) // FIXME_NVHPC
if (!std::is_same_v<ExecutionSpace, Kokkos::Cuda>)
#endif
test_issue_1160_sort<ExecutionSpace>();
test_issue_4978_sort<ExecutionSpace>();
test_sort_integer_overflow<ExecutionSpace, long long>();
test_sort_integer_overflow<ExecutionSpace, unsigned long long>();
test_sort_integer_overflow<ExecutionSpace, int>();
SortImpl::test_issue_4978_impl<ExecutionSpace>();
}
} // namespace Impl
TEST(TEST_CATEGORY, SortEmptyView) {
using ExecutionSpace = TEST_EXECSPACE;
// does not matter if we use int or something else
Kokkos::View<int*, ExecutionSpace> v("v", 0);
// TODO check the synchronous behavior of the calls below
ASSERT_NO_THROW(Kokkos::sort(ExecutionSpace(), v));
ASSERT_NO_THROW(Kokkos::sort(v));
}
} // namespace Test
#endif /* KOKKOS_ALGORITHMS_UNITTESTS_TESTSORT_HPP */
#endif

27
lib/kokkos/algorithms/unit_tests/TestSortCommon.hpp
View File

@ -1,27 +0,0 @@
//@HEADER
// ************************************************************************
//
// Kokkos v. 4.0
// Copyright (2022) National Technology & Engineering
// Solutions of Sandia, LLC (NTESS).
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// the U.S. Government retains certain rights in this software.
//
// Part of Kokkos, under the Apache License v2.0 with LLVM Exceptions.
// See https://kokkos.org/LICENSE for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//@HEADER
#ifndef KOKKOS_ALGORITHMS_UNITTESTS_TESTSORT_COMMON_HPP
#define KOKKOS_ALGORITHMS_UNITTESTS_TESTSORT_COMMON_HPP
#include <TestSort.hpp>
namespace Test {
TEST(TEST_CATEGORY, SortUnsigned) {
Impl::test_sort<TEST_EXECSPACE, unsigned>(171);
}
} // namespace Test
#endif

10
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsAdjacentDifference.cpp
View File

@ -157,7 +157,7 @@ void verify_data(TestViewType test_view, GoldViewType gold) {
const auto gold_h = create_mirror_view_and_copy(Kokkos::HostSpace(), gold);
for (std::size_t i = 0; i < test_view.extent(0); ++i) {
EXPECT_EQ(gold_h(i), test_view_dc_h(i));
ASSERT_EQ(gold_h(i), test_view_dc_h(i));
}
}
@ -197,7 +197,7 @@ void run_single_scenario(const InfoType& scenario_info,
auto res1 = KE::adjacent_difference(exespace(), KE::cbegin(view_from),
KE::cend(view_from),
KE::begin(view_dest), args...);
EXPECT_EQ(res1, KE::end(view_dest));
ASSERT_EQ(res1, KE::end(view_dest));
verify_data(view_dest, gold);
}
@ -207,7 +207,7 @@ void run_single_scenario(const InfoType& scenario_info,
auto res2 = KE::adjacent_difference(
"label", exespace(), KE::cbegin(view_from), KE::cend(view_from),
KE::begin(view_dest), args...);
EXPECT_EQ(res2, KE::end(view_dest));
ASSERT_EQ(res2, KE::end(view_dest));
verify_data(view_dest, gold);
}
@ -216,7 +216,7 @@ void run_single_scenario(const InfoType& scenario_info,
create_view<ValueType>(Tag{}, view_ext, "adj_diff_dest_view");
auto res3 =
KE::adjacent_difference(exespace(), view_from, view_dest, args...);
EXPECT_EQ(res3, KE::end(view_dest));
ASSERT_EQ(res3, KE::end(view_dest));
verify_data(view_dest, gold);
}
@ -225,7 +225,7 @@ void run_single_scenario(const InfoType& scenario_info,
create_view<ValueType>(Tag{}, view_ext, "adj_diff_dest_view");
auto res4 = KE::adjacent_difference("label", exespace(), view_from,
view_dest, args...);
EXPECT_EQ(res4, KE::end(view_dest));
ASSERT_EQ(res4, KE::end(view_dest));
verify_data(view_dest, gold);
}

8
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsAdjacentFind.cpp
View File

@ -229,7 +229,7 @@ void verify(DiffType my_diff, ViewType view, Args... args) {
my_std_adjacent_find(KE::cbegin(view_h), KE::cend(view_h), args...);
const auto std_diff = std_r - KE::cbegin(view_h);
EXPECT_EQ(my_diff, std_diff);
ASSERT_EQ(my_diff, std_diff);
}
template <class Tag, class ValueType, class InfoType, class... Args>
@ -287,12 +287,6 @@ void run_all_scenarios() {
}
TEST(std_algorithms_nonmod_seq_ops, adjacent_find) {
#if defined(KOKKOS_ENABLE_CUDA) && \
defined(KOKKOS_COMPILER_NVHPC) // FIXME_NVHPC
if constexpr (std::is_same_v<exespace, Kokkos::Cuda>) {
GTEST_SKIP() << "FIXME wrong result";
}
#endif
run_all_scenarios<DynamicTag, int>();
run_all_scenarios<DynamicTag, double>();
run_all_scenarios<StridedThreeTag, int>();

6
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsAllAnyNoneOf.cpp
View File

@ -147,12 +147,6 @@ void run_all_scenarios() {
}
TEST(std_algorithms_all_any_none_of_test, test) {
#if defined(KOKKOS_ENABLE_CUDA) && \
defined(KOKKOS_COMPILER_NVHPC) // FIXME_NVHPC
if constexpr (std::is_same_v<exespace, Kokkos::Cuda>) {
GTEST_SKIP() << "FIXME wrong result";
}
#endif
run_all_scenarios<DynamicTag, double>();
run_all_scenarios<StridedTwoTag, int>();
run_all_scenarios<StridedThreeTag, unsigned>();

8
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsCommon.hpp
View File

@ -110,7 +110,7 @@ verify_values(ValueType expected, const ViewType view) {
"Non-matching value types of view and reference value");
auto view_h = Kokkos::create_mirror_view_and_copy(Kokkos::HostSpace(), view);
for (std::size_t i = 0; i < view_h.extent(0); i++) {
EXPECT_EQ(expected, view_h(i));
ASSERT_EQ(expected, view_h(i));
}
}
@ -130,7 +130,7 @@ verify_values(ValueType expected, const ViewType view) {
auto view_h =
Kokkos::create_mirror_view_and_copy(Kokkos::HostSpace(), tmpView);
for (std::size_t i = 0; i < view_h.extent(0); i++) {
EXPECT_EQ(expected, view_h(i));
ASSERT_EQ(expected, view_h(i));
}
}
@ -147,7 +147,7 @@ compare_views(ViewType1 expected, const ViewType2 actual) {
Kokkos::create_mirror_view_and_copy(Kokkos::HostSpace(), actual);
for (std::size_t i = 0; i < expected_h.extent(0); i++) {
EXPECT_EQ(expected_h(i), actual_h(i));
ASSERT_EQ(expected_h(i), actual_h(i));
}
}
@ -171,7 +171,7 @@ compare_views(ViewType1 expected, const ViewType2 actual) {
Kokkos::create_mirror_view_and_copy(Kokkos::HostSpace(), expected);
for (std::size_t i = 0; i < expected_h.extent(0); i++) {
EXPECT_EQ(expected_h(i), actual_h(i));
ASSERT_EQ(expected_h(i), actual_h(i));
}
}

4
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsConstraints.cpp
View File

@ -42,8 +42,8 @@ TEST(std_algorithms, is_admissible_to_std_algorithms) {
using strided_view_1d_t = Kokkos::View<value_type*, Kokkos::LayoutStride>;
Kokkos::LayoutStride layout1d{extent0, 2};
strided_view_1d_t strided_view_1d{"std-algo-test-1d-strided-view", layout1d};
EXPECT_EQ(layout1d.dimension[0], 13u);
EXPECT_EQ(layout1d.stride[0], 2u);
ASSERT_EQ(layout1d.dimension[0], 13u);
ASSERT_EQ(layout1d.stride[0], 2u);
// they are admissible
KE::Impl::static_assert_is_admissible_to_kokkos_std_algorithms(
static_view_1d);

68
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsCopyIf.cpp
View File

@ -135,49 +135,49 @@ void verify_data(const std::string& name, ViewTypeFrom view_from,
}
else if (name == "one-element-a") {
EXPECT_EQ(view_test_h(0), static_cast<value_type>(0));
ASSERT_EQ(view_test_h(0), static_cast<value_type>(0));
}
else if (name == "one-element-b") {
EXPECT_EQ(view_test_h(0), static_cast<value_type>(2));
ASSERT_EQ(view_test_h(0), static_cast<value_type>(2));
}
else if (name == "two-elements-a") {
EXPECT_EQ(view_test_h(0), static_cast<value_type>(2));
EXPECT_EQ(view_test_h(1), static_cast<value_type>(0));
ASSERT_EQ(view_test_h(0), static_cast<value_type>(2));
ASSERT_EQ(view_test_h(1), static_cast<value_type>(0));
}
else if (name == "two-elements-b") {
EXPECT_EQ(view_test_h(0), static_cast<value_type>(2));
EXPECT_EQ(view_test_h(1), static_cast<value_type>(0));
ASSERT_EQ(view_test_h(0), static_cast<value_type>(2));
ASSERT_EQ(view_test_h(1), static_cast<value_type>(0));
}
else if (name == "small-a") {
EXPECT_EQ(view_test_h(0), static_cast<value_type>(-4));
EXPECT_EQ(view_test_h(1), static_cast<value_type>(-2));
EXPECT_EQ(view_test_h(2), static_cast<value_type>(0));
EXPECT_EQ(view_test_h(3), static_cast<value_type>(2));
EXPECT_EQ(view_test_h(4), static_cast<value_type>(4));
EXPECT_EQ(view_test_h(5), static_cast<value_type>(0));
EXPECT_EQ(view_test_h(6), static_cast<value_type>(0));
EXPECT_EQ(view_test_h(7), static_cast<value_type>(0));
EXPECT_EQ(view_test_h(8), static_cast<value_type>(0));
ASSERT_EQ(view_test_h(0), static_cast<value_type>(-4));
ASSERT_EQ(view_test_h(1), static_cast<value_type>(-2));
ASSERT_EQ(view_test_h(2), static_cast<value_type>(0));
ASSERT_EQ(view_test_h(3), static_cast<value_type>(2));
ASSERT_EQ(view_test_h(4), static_cast<value_type>(4));
ASSERT_EQ(view_test_h(5), static_cast<value_type>(0));
ASSERT_EQ(view_test_h(6), static_cast<value_type>(0));
ASSERT_EQ(view_test_h(7), static_cast<value_type>(0));
ASSERT_EQ(view_test_h(8), static_cast<value_type>(0));
}
else if (name == "small-b") {
EXPECT_EQ(view_test_h(0), static_cast<value_type>(22));
EXPECT_EQ(view_test_h(1), static_cast<value_type>(-12));
EXPECT_EQ(view_test_h(2), static_cast<value_type>(22));
EXPECT_EQ(view_test_h(3), static_cast<value_type>(-12));
EXPECT_EQ(view_test_h(4), static_cast<value_type>(22));
EXPECT_EQ(view_test_h(5), static_cast<value_type>(-12));
EXPECT_EQ(view_test_h(6), static_cast<value_type>(22));
EXPECT_EQ(view_test_h(7), static_cast<value_type>(-12));
EXPECT_EQ(view_test_h(8), static_cast<value_type>(22));
EXPECT_EQ(view_test_h(9), static_cast<value_type>(-12));
EXPECT_EQ(view_test_h(10), static_cast<value_type>(22));
EXPECT_EQ(view_test_h(11), static_cast<value_type>(-12));
EXPECT_EQ(view_test_h(12), static_cast<value_type>(22));
ASSERT_EQ(view_test_h(0), static_cast<value_type>(22));
ASSERT_EQ(view_test_h(1), static_cast<value_type>(-12));
ASSERT_EQ(view_test_h(2), static_cast<value_type>(22));
ASSERT_EQ(view_test_h(3), static_cast<value_type>(-12));
ASSERT_EQ(view_test_h(4), static_cast<value_type>(22));
ASSERT_EQ(view_test_h(5), static_cast<value_type>(-12));
ASSERT_EQ(view_test_h(6), static_cast<value_type>(22));
ASSERT_EQ(view_test_h(7), static_cast<value_type>(-12));
ASSERT_EQ(view_test_h(8), static_cast<value_type>(22));
ASSERT_EQ(view_test_h(9), static_cast<value_type>(-12));
ASSERT_EQ(view_test_h(10), static_cast<value_type>(22));
ASSERT_EQ(view_test_h(11), static_cast<value_type>(-12));
ASSERT_EQ(view_test_h(12), static_cast<value_type>(22));
}
else if (name == "medium" || name == "large") {
@ -190,14 +190,14 @@ void verify_data(const std::string& name, ViewTypeFrom view_from,
std::size_t count = 0;
for (std::size_t i = 0; i < view_from_h.extent(0); ++i) {
if (pred(view_from_h(i))) {
EXPECT_EQ(view_test_h(count), view_from_h(i));
ASSERT_EQ(view_test_h(count), view_from_h(i));
count++;
}
}
// all other entries of test view should be zero
for (; count < view_test_h.extent(0); ++count) {
// std::cout << count << '\n';
EXPECT_EQ(view_test_h(count), value_type(0));
ASSERT_EQ(view_test_h(count), value_type(0));
}
}
@ -226,7 +226,7 @@ void run_single_scenario(const InfoType& scenario_info) {
auto rit = KE::copy_if(exespace(), KE::cbegin(view_from),
KE::cend(view_from), KE::begin(view_dest), pred);
verify_data(name, view_from, view_dest, pred);
EXPECT_EQ(rit, (KE::begin(view_dest) + n));
ASSERT_EQ(rit, (KE::begin(view_dest) + n));
}
{
@ -235,7 +235,7 @@ void run_single_scenario(const InfoType& scenario_info) {
auto rit = KE::copy_if("label", exespace(), KE::cbegin(view_from),
KE::cend(view_from), KE::begin(view_dest), pred);
verify_data(name, view_from, view_dest, pred);
EXPECT_EQ(rit, (KE::begin(view_dest) + n));
ASSERT_EQ(rit, (KE::begin(view_dest) + n));
}
{
@ -243,7 +243,7 @@ void run_single_scenario(const InfoType& scenario_info) {
auto view_dest = create_view<ValueType>(Tag{}, view_ext, "copy_if_dest");
auto rit = KE::copy_if(exespace(), view_from, view_dest, pred);
verify_data(name, view_from, view_dest, pred);
EXPECT_EQ(rit, (KE::begin(view_dest) + n));
ASSERT_EQ(rit, (KE::begin(view_dest) + n));
}
{
@ -251,7 +251,7 @@ void run_single_scenario(const InfoType& scenario_info) {
auto view_dest = create_view<ValueType>(Tag{}, view_ext, "copy_if_dest");
auto rit = KE::copy_if("label", exespace(), view_from, view_dest, pred);
verify_data(name, view_from, view_dest, pred);
EXPECT_EQ(rit, (KE::begin(view_dest) + n));
ASSERT_EQ(rit, (KE::begin(view_dest) + n));
}
Kokkos::fence();

20
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsCount.cpp
View File

@ -35,13 +35,13 @@ void test_count(const ViewType view) {
const value_t count_value = 0;
const auto std_result =
std::count(KE::cbegin(expected), KE::cend(expected), count_value);
EXPECT_EQ(view.extent(0), size_t(std_result));
ASSERT_EQ(view.extent(0), size_t(std_result));
// pass const iterators
EXPECT_EQ(std_result, KE::count(exespace(), KE::cbegin(view),
ASSERT_EQ(std_result, KE::count(exespace(), KE::cbegin(view),
KE::cend(view), count_value));
// pass view
EXPECT_EQ(std_result, KE::count(exespace(), view, count_value));
ASSERT_EQ(std_result, KE::count(exespace(), view, count_value));
}
{
@ -50,10 +50,10 @@ void test_count(const ViewType view) {
std::count(KE::cbegin(expected), KE::cend(expected), count_value);
// pass iterators
EXPECT_EQ(std_result, KE::count("label", exespace(), KE::begin(view),
ASSERT_EQ(std_result, KE::count("label", exespace(), KE::begin(view),
KE::end(view), count_value));
// pass view
EXPECT_EQ(std_result, KE::count("label", exespace(), view, count_value));
ASSERT_EQ(std_result, KE::count("label", exespace(), view, count_value));
}
}
@ -67,24 +67,24 @@ void test_count_if(const ViewType view) {
// no positive elements (all zeroes)
const auto predicate = IsPositiveFunctor<value_type>();
EXPECT_EQ(0,
ASSERT_EQ(0,
std::count_if(KE::begin(expected), KE::end(expected), predicate));
// pass iterators
EXPECT_EQ(
ASSERT_EQ(
0, KE::count_if(exespace(), KE::begin(view), KE::end(view), predicate));
// pass view
EXPECT_EQ(0, KE::count_if(exespace(), view, predicate));
ASSERT_EQ(0, KE::count_if(exespace(), view, predicate));
fill_views_inc(view, expected);
const auto std_result =
std::count_if(KE::begin(expected), KE::end(expected), predicate);
// pass const iterators
EXPECT_EQ(std_result, KE::count_if("label", exespace(), KE::cbegin(view),
ASSERT_EQ(std_result, KE::count_if("label", exespace(), KE::cbegin(view),
KE::cend(view), predicate));
// pass view
EXPECT_EQ(std_result, KE::count_if("label", exespace(), view, predicate));
ASSERT_EQ(std_result, KE::count_if("label", exespace(), view, predicate));
}
template <class Tag, class ValueType>

58
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsExclusiveScan.cpp
View File

@ -157,7 +157,7 @@ void verify_data(ViewType1 data_view, // contains data
// << gold_h(i) << " " << test_view_h(i) << " "
// << std::abs(gold_h(i) - test_view_h(i)) << std::endl;
if (std::is_same<gold_view_value_type, int>::value) {
EXPECT_EQ(gold_h(i), test_view_h(i));
ASSERT_EQ(gold_h(i), test_view_h(i));
} else {
const auto error =
std::abs(static_cast<double>(gold_h(i) - test_view_h(i)));
@ -213,7 +213,7 @@ void run_single_scenario_default_op(const InfoType& scenario_info,
auto r = KE::exclusive_scan(exespace(), KE::cbegin(view_from),
KE::cend(view_from), KE::begin(view_dest),
init_value);
EXPECT_EQ(r, KE::end(view_dest));
ASSERT_EQ(r, KE::end(view_dest));
verify_data(view_from, view_dest, init_value, default_op());
}
@ -222,14 +222,14 @@ void run_single_scenario_default_op(const InfoType& scenario_info,
auto r = KE::exclusive_scan("label", exespace(), KE::cbegin(view_from),
KE::cend(view_from), KE::begin(view_dest),
init_value);
EXPECT_EQ(r, KE::end(view_dest));
ASSERT_EQ(r, KE::end(view_dest));
verify_data(view_from, view_dest, init_value, default_op());
}
{
fill_zero(view_dest);
auto r = KE::exclusive_scan(exespace(), view_from, view_dest, init_value);
EXPECT_EQ(r, KE::end(view_dest));
ASSERT_EQ(r, KE::end(view_dest));
verify_data(view_from, view_dest, init_value, default_op());
}
@ -237,7 +237,7 @@ void run_single_scenario_default_op(const InfoType& scenario_info,
fill_zero(view_dest);
auto r = KE::exclusive_scan("label", exespace(), view_from, view_dest,
init_value);
EXPECT_EQ(r, KE::end(view_dest));
ASSERT_EQ(r, KE::end(view_dest));
verify_data(view_from, view_dest, init_value, default_op());
}
@ -263,7 +263,7 @@ void run_single_scenario_custom_op(const InfoType& scenario_info,
auto r = KE::exclusive_scan(exespace(), KE::cbegin(view_from),
KE::cend(view_from), KE::begin(view_dest),
init_value, bop);
EXPECT_EQ(r, KE::end(view_dest));
ASSERT_EQ(r, KE::end(view_dest));
verify_data(view_from, view_dest, init_value, bop);
}
@ -272,7 +272,7 @@ void run_single_scenario_custom_op(const InfoType& scenario_info,
auto r = KE::exclusive_scan("label", exespace(), KE::cbegin(view_from),
KE::cend(view_from), KE::begin(view_dest),
init_value, bop);
EXPECT_EQ(r, KE::end(view_dest));
ASSERT_EQ(r, KE::end(view_dest));
verify_data(view_from, view_dest, init_value, bop);
}
@ -280,7 +280,7 @@ void run_single_scenario_custom_op(const InfoType& scenario_info,
fill_zero(view_dest);
auto r =
KE::exclusive_scan(exespace(), view_from, view_dest, init_value, bop);
EXPECT_EQ(r, KE::end(view_dest));
ASSERT_EQ(r, KE::end(view_dest));
verify_data(view_from, view_dest, init_value, bop);
}
@ -288,7 +288,7 @@ void run_single_scenario_custom_op(const InfoType& scenario_info,
fill_zero(view_dest);
auto r = KE::exclusive_scan("label", exespace(), view_from, view_dest,
init_value, bop);
EXPECT_EQ(r, KE::end(view_dest));
ASSERT_EQ(r, KE::end(view_dest));
verify_data(view_from, view_dest, init_value, bop);
}
@ -344,6 +344,46 @@ TEST(std_algorithms_numeric_ops_test, exclusive_scan) {
run_exclusive_scan_all_scenarios<StridedThreeTag, CustomValueType>();
}
TEST(std_algorithms_numeric_ops_test, exclusive_scan_functor) {
int dummy = 0;
using view_type = Kokkos::View<int*, exespace>;
view_type dummy_view("dummy_view", 0);
using functor_type = Kokkos::Experimental::Impl::ExclusiveScanDefaultFunctor<
exespace, int, int, view_type, view_type>;
functor_type functor(dummy, dummy_view, dummy_view);
using value_type = functor_type::value_type;
value_type value1;
functor.init(value1);
ASSERT_EQ(value1.val, 0);
ASSERT_EQ(value1.is_initial, true);
value_type value2;
value2.val = 1;
value2.is_initial = false;
functor.join(value1, value2);
ASSERT_EQ(value1.val, 1);
ASSERT_EQ(value1.is_initial, false);
functor.init(value1);
functor.join(value2, value1);
ASSERT_EQ(value2.val, 1);
ASSERT_EQ(value2.is_initial, false);
functor.init(value2);
functor.join(value2, value1);
ASSERT_EQ(value2.val, 0);
ASSERT_EQ(value2.is_initial, true);
value1.val = 1;
value1.is_initial = false;
value2.val = 2;
value2.is_initial = false;
functor.join(value2, value1);
ASSERT_EQ(value2.val, 3);
ASSERT_EQ(value2.is_initial, false);
}
} // namespace EScan
} // namespace stdalgos
} // namespace Test

36
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsFind.cpp
View File

@ -34,14 +34,14 @@ void test_find(const ViewType view) {
constexpr value_t find_value = 13;
// value not found, return last
EXPECT_EQ(KE::end(expected),
ASSERT_EQ(KE::end(expected),
std::find(KE::begin(expected), KE::end(expected), find_value));
// pass const iterators, returns const iterator
EXPECT_EQ(KE::cend(view),
ASSERT_EQ(KE::cend(view),
KE::find(exespace(), KE::cbegin(view), KE::cend(view), find_value));
// pass view, returns iterator
EXPECT_EQ(KE::end(view), KE::find(exespace(), view, find_value));
ASSERT_EQ(KE::end(view), KE::find(exespace(), view, find_value));
fill_views_inc(view, expected);
@ -50,10 +50,10 @@ void test_find(const ViewType view) {
auto distance = std::distance(KE::begin(expected), std_result);
// pass iterators, returns iterator
EXPECT_EQ(KE::begin(view) + distance,
ASSERT_EQ(KE::begin(view) + distance,
KE::find(exespace(), KE::begin(view), KE::end(view), find_value));
// pass view, returns iterator
EXPECT_EQ(KE::begin(view) + distance, KE::find(exespace(), view, find_value));
ASSERT_EQ(KE::begin(view) + distance, KE::find(exespace(), view, find_value));
}
template <class ViewType>
@ -67,15 +67,15 @@ void test_find_if(const ViewType view) {
const auto not_equals_zero = NotEqualsZeroFunctor<value_type>();
// value not found, return last
EXPECT_EQ(
ASSERT_EQ(
KE::end(expected),
std::find_if(KE::begin(expected), KE::end(expected), not_equals_zero));
// pass iterators, returns iterator
EXPECT_EQ(KE::end(view), KE::find_if(exespace(), KE::begin(view),
ASSERT_EQ(KE::end(view), KE::find_if(exespace(), KE::begin(view),
KE::end(view), not_equals_zero));
// pass view, returns iterator
EXPECT_EQ(KE::end(view), KE::find_if(exespace(), view, not_equals_zero));
ASSERT_EQ(KE::end(view), KE::find_if(exespace(), view, not_equals_zero));
fill_views_inc(view, expected);
@ -86,11 +86,11 @@ void test_find_if(const ViewType view) {
auto distance = std::distance(KE::begin(expected), std_result);
// pass const iterators, returns const iterator
EXPECT_EQ(
ASSERT_EQ(
KE::cbegin(view) + distance,
KE::find_if(exespace(), KE::cbegin(view), KE::cend(view), equals_val));
// pass view, returns iterator
EXPECT_EQ(KE::begin(view) + distance,
ASSERT_EQ(KE::begin(view) + distance,
KE::find_if(exespace(), view, equals_val));
}
@ -105,15 +105,15 @@ void test_find_if_not(const ViewType view) {
const auto not_equals_zero = NotEqualsZeroFunctor<value_type>();
// first value matches
EXPECT_EQ(KE::begin(expected),
ASSERT_EQ(KE::begin(expected),
std::find_if_not(KE::begin(expected), KE::end(expected),
not_equals_zero));
// pass iterators, returns iterator
EXPECT_EQ(KE::begin(view), KE::find_if_not(exespace(), KE::begin(view),
ASSERT_EQ(KE::begin(view), KE::find_if_not(exespace(), KE::begin(view),
KE::end(view), not_equals_zero));
// pass view, returns iterator
EXPECT_EQ(KE::begin(view),
ASSERT_EQ(KE::begin(view),
KE::find_if_not(exespace(), view, not_equals_zero));
fill_views_inc(view, expected);
@ -124,11 +124,11 @@ void test_find_if_not(const ViewType view) {
auto distance = std::distance(KE::begin(expected), std_result);
// pass const iterators, returns const iterator
EXPECT_EQ(KE::cbegin(view) + distance,
ASSERT_EQ(KE::cbegin(view) + distance,
KE::find_if_not(exespace(), KE::cbegin(view), KE::cend(view),
equals_zero));
// pass view, returns const iterator
EXPECT_EQ(KE::begin(view) + distance,
ASSERT_EQ(KE::begin(view) + distance,
KE::find_if_not(exespace(), view, equals_zero));
}
@ -151,12 +151,6 @@ void run_all_scenarios() {
}
TEST(std_algorithms_find_test, test) {
#if defined(KOKKOS_ENABLE_CUDA) && \
defined(KOKKOS_COMPILER_NVHPC) // FIXME_NVHPC
if constexpr (std::is_same_v<exespace, Kokkos::Cuda>) {
GTEST_SKIP() << "FIXME wrong result";
}
#endif
run_all_scenarios<DynamicTag, double>();
run_all_scenarios<StridedTwoTag, int>();
run_all_scenarios<StridedThreeTag, unsigned>();

14
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsFindEnd.cpp
View File

@ -282,7 +282,7 @@ void run_single_scenario(const InfoType& scenario_info, std::size_t seq_ext,
const auto mydiff = myrit - KE::cbegin(view);
const auto stddiff = stdrit - KE::cbegin(view_h);
// std::cout << "result : " << mydiff << " " << stddiff << std::endl;
EXPECT_EQ(mydiff, stddiff);
ASSERT_EQ(mydiff, stddiff);
}
{
@ -291,21 +291,21 @@ void run_single_scenario(const InfoType& scenario_info, std::size_t seq_ext,
KE::cbegin(s_view), KE::cend(s_view), args...);
const auto mydiff = myrit - KE::cbegin(view);
const auto stddiff = stdrit - KE::cbegin(view_h);
EXPECT_EQ(mydiff, stddiff);
ASSERT_EQ(mydiff, stddiff);
}
{
auto myrit = KE::find_end(exespace(), view, s_view, args...);
const auto mydiff = myrit - KE::begin(view);
const auto stddiff = stdrit - KE::cbegin(view_h);
EXPECT_EQ(mydiff, stddiff);
ASSERT_EQ(mydiff, stddiff);
}
{
auto myrit = KE::find_end("label", exespace(), view, s_view, args...);
const auto mydiff = myrit - KE::begin(view);
const auto stddiff = stdrit - KE::cbegin(view_h);
EXPECT_EQ(mydiff, stddiff);
ASSERT_EQ(mydiff, stddiff);
}
Kokkos::fence();
@ -348,12 +348,6 @@ void run_all_scenarios() {
}
TEST(std_algorithms_non_mod_seq_ops, find_end) {
#if defined(KOKKOS_ENABLE_CUDA) && \
defined(KOKKOS_COMPILER_NVHPC) // FIXME_NVHPC
if constexpr (std::is_same_v<exespace, Kokkos::Cuda>) {
GTEST_SKIP() << "FIXME wrong result";
}
#endif
run_all_scenarios<DynamicTag, int>();
run_all_scenarios<StridedThreeTag, int>();
}

14
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsFindFirstOf.cpp
View File

@ -201,7 +201,7 @@ void run_single_scenario(const InfoType& scenario_info, std::size_t seq_ext,
KE::cbegin(s_view), KE::cend(s_view), args...);
const auto mydiff = myrit - KE::cbegin(view);
const auto stddiff = stdrit - KE::cbegin(view_h);
EXPECT_EQ(mydiff, stddiff);
ASSERT_EQ(mydiff, stddiff);
}
{
@ -210,21 +210,21 @@ void run_single_scenario(const InfoType& scenario_info, std::size_t seq_ext,
KE::cbegin(s_view), KE::cend(s_view), args...);
const auto mydiff = myrit - KE::cbegin(view);
const auto stddiff = stdrit - KE::cbegin(view_h);
EXPECT_EQ(mydiff, stddiff);
ASSERT_EQ(mydiff, stddiff);
}
{
auto myrit = KE::find_first_of(exespace(), view, s_view, args...);
const auto mydiff = myrit - KE::begin(view);
const auto stddiff = stdrit - KE::cbegin(view_h);
EXPECT_EQ(mydiff, stddiff);
ASSERT_EQ(mydiff, stddiff);
}
{
auto myrit = KE::find_first_of("label", exespace(), view, s_view, args...);
const auto mydiff = myrit - KE::begin(view);
const auto stddiff = stdrit - KE::cbegin(view_h);
EXPECT_EQ(mydiff, stddiff);
ASSERT_EQ(mydiff, stddiff);
}
Kokkos::fence();
@ -264,12 +264,6 @@ void run_all_scenarios() {
}
TEST(std_algorithms_non_mod_seq_ops, find_first_of) {
#if defined(KOKKOS_ENABLE_CUDA) && \
defined(KOKKOS_COMPILER_NVHPC) // FIXME_NVHPC
if constexpr (std::is_same_v<exespace, Kokkos::Cuda>) {
GTEST_SKIP() << "FIXME wrong result";
}
#endif
run_all_scenarios<DynamicTag, int>();
run_all_scenarios<StridedThreeTag, int>();
}

8
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsForEach.cpp
View File

@ -91,23 +91,23 @@ void test_for_each_n(const ViewType view) {
const auto non_mod_functor = NoOpNonMutableFunctor<value_t>();
// pass const iterators, functor takes const ref
EXPECT_EQ(KE::cbegin(view) + n,
ASSERT_EQ(KE::cbegin(view) + n,
KE::for_each_n(exespace(), KE::cbegin(view), n, non_mod_functor));
verify_values(value_t{0}, view);
// pass view, functor takes const ref
EXPECT_EQ(KE::begin(view) + n,
ASSERT_EQ(KE::begin(view) + n,
KE::for_each_n(exespace(), view, n, non_mod_functor));
verify_values(value_t{0}, view);
// pass iterators, functor takes non-const ref
const auto mod_functor = IncrementElementWiseFunctor<value_t>();
EXPECT_EQ(KE::begin(view) + n,
ASSERT_EQ(KE::begin(view) + n,
KE::for_each_n(exespace(), KE::begin(view), n, mod_functor));
verify_values(value_t{1}, view);
// pass view, functor takes non-const ref
EXPECT_EQ(KE::begin(view) + n,
ASSERT_EQ(KE::begin(view) + n,
KE::for_each_n("label", exespace(), view, n, mod_functor));
verify_values(value_t{2}, view);
}

57
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsInclusiveScan.cpp
View File

@ -171,7 +171,7 @@ void verify_data(ViewType1 data_view, // contains data
// << std::abs(gold_h(i) - test_view_h(i)) << std::endl;
if (std::is_same<gold_view_value_type, int>::value) {
EXPECT_EQ(gold_h(i), test_view_h(i));
ASSERT_EQ(gold_h(i), test_view_h(i));
} else {
const auto error =
std::abs(static_cast<double>(gold_h(i) - test_view_h(i)));
@ -224,7 +224,7 @@ void run_single_scenario_default_op(const InfoType& scenario_info) {
fill_zero(view_dest);
auto r = KE::inclusive_scan(exespace(), KE::cbegin(view_from),
KE::cend(view_from), KE::begin(view_dest));
EXPECT_EQ(r, KE::end(view_dest));
ASSERT_EQ(r, KE::end(view_dest));
verify_data(view_from, view_dest, default_op());
}
@ -232,21 +232,21 @@ void run_single_scenario_default_op(const InfoType& scenario_info) {
fill_zero(view_dest);
auto r = KE::inclusive_scan("label", exespace(), KE::cbegin(view_from),
KE::cend(view_from), KE::begin(view_dest));
EXPECT_EQ(r, KE::end(view_dest));
ASSERT_EQ(r, KE::end(view_dest));
verify_data(view_from, view_dest, default_op());
}
{
fill_zero(view_dest);
auto r = KE::inclusive_scan(exespace(), view_from, view_dest);
EXPECT_EQ(r, KE::end(view_dest));
ASSERT_EQ(r, KE::end(view_dest));
verify_data(view_from, view_dest, default_op());
}
{
fill_zero(view_dest);
auto r = KE::inclusive_scan("label", exespace(), view_from, view_dest);
EXPECT_EQ(r, KE::end(view_dest));
ASSERT_EQ(r, KE::end(view_dest));
verify_data(view_from, view_dest, default_op());
}
@ -279,7 +279,7 @@ void run_single_scenario_custom_op(const InfoType& scenario_info, BinaryOp bop,
auto r = KE::inclusive_scan(exespace(), KE::cbegin(view_from),
KE::cend(view_from), KE::begin(view_dest), bop,
args...);
EXPECT_EQ(r, KE::end(view_dest));
ASSERT_EQ(r, KE::end(view_dest));
verify_data(view_from, view_dest, bop, args...);
}
@ -288,14 +288,14 @@ void run_single_scenario_custom_op(const InfoType& scenario_info, BinaryOp bop,
auto r = KE::inclusive_scan("label", exespace(), KE::cbegin(view_from),
KE::cend(view_from), KE::begin(view_dest), bop,
args...);
EXPECT_EQ(r, KE::end(view_dest));
ASSERT_EQ(r, KE::end(view_dest));
verify_data(view_from, view_dest, bop, args...);
}
{
fill_zero(view_dest);
auto r = KE::inclusive_scan(exespace(), view_from, view_dest, bop, args...);
EXPECT_EQ(r, KE::end(view_dest));
ASSERT_EQ(r, KE::end(view_dest));
verify_data(view_from, view_dest, bop, args...);
}
@ -303,7 +303,7 @@ void run_single_scenario_custom_op(const InfoType& scenario_info, BinaryOp bop,
fill_zero(view_dest);
auto r = KE::inclusive_scan("label", exespace(), view_from, view_dest, bop,
args...);
EXPECT_EQ(r, KE::end(view_dest));
ASSERT_EQ(r, KE::end(view_dest));
verify_data(view_from, view_dest, bop, args...);
}
@ -353,6 +353,45 @@ TEST(std_algorithms_numeric_ops_test, inclusive_scan) {
run_inclusive_scan_all_scenarios<StridedThreeTag, CustomValueType>();
}
TEST(std_algorithms_numeric_ops_test, inclusive_scan_functor) {
using view_type = Kokkos::View<int*, exespace>;
view_type dummy_view("dummy_view", 0);
using functor_type = Kokkos::Experimental::Impl::InclusiveScanDefaultFunctor<
exespace, int, int, view_type, view_type>;
functor_type functor(dummy_view, dummy_view);
using value_type = functor_type::value_type;
value_type value1;
functor.init(value1);
ASSERT_EQ(value1.val, 0);
ASSERT_EQ(value1.is_initial, true);
value_type value2;
value2.val = 1;
value2.is_initial = false;
functor.join(value1, value2);
ASSERT_EQ(value1.val, 1);
ASSERT_EQ(value1.is_initial, false);
functor.init(value1);
functor.join(value2, value1);
ASSERT_EQ(value2.val, 1);
ASSERT_EQ(value2.is_initial, false);
functor.init(value2);
functor.join(value2, value1);
ASSERT_EQ(value2.val, 0);
ASSERT_EQ(value2.is_initial, true);
value1.val = 1;
value1.is_initial = false;
value2.val = 2;
value2.is_initial = false;
functor.join(value2, value1);
ASSERT_EQ(value2.val, 3);
ASSERT_EQ(value2.is_initial, false);
}
} // namespace IncScan
} // namespace stdalgos
} // namespace Test

22
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsIsSortedUntil.cpp
View File

@ -145,10 +145,10 @@ void run_single_scenario(const InfoType& scenario_info) {
KE::is_sorted_until("label", exespace(), KE::begin(view), KE::end(view));
auto r3 = KE::is_sorted_until(exespace(), view);
auto r4 = KE::is_sorted_until("label", exespace(), view);
EXPECT_EQ(r1, gold);
EXPECT_EQ(r2, gold);
EXPECT_EQ(r3, gold);
EXPECT_EQ(r4, gold);
ASSERT_EQ(r1, gold);
ASSERT_EQ(r2, gold);
ASSERT_EQ(r3, gold);
ASSERT_EQ(r4, gold);
#if !defined KOKKOS_ENABLE_OPENMPTARGET
CustomLessThanComparator<ValueType, ValueType> comp;
@ -160,10 +160,10 @@ void run_single_scenario(const InfoType& scenario_info) {
auto r8 = KE::is_sorted_until("label", exespace(), view, comp);
#endif
EXPECT_EQ(r1, gold);
EXPECT_EQ(r2, gold);
EXPECT_EQ(r3, gold);
EXPECT_EQ(r4, gold);
ASSERT_EQ(r1, gold);
ASSERT_EQ(r2, gold);
ASSERT_EQ(r3, gold);
ASSERT_EQ(r4, gold);
Kokkos::fence();
}
@ -185,12 +185,6 @@ void run_is_sorted_until_all_scenarios() {
}
TEST(std_algorithms_sorting_ops_test, is_sorted_until) {
#if defined(KOKKOS_ENABLE_CUDA) && \
defined(KOKKOS_COMPILER_NVHPC) // FIXME_NVHPC
if constexpr (std::is_same_v<exespace, Kokkos::Cuda>) {
GTEST_SKIP() << "FIXME wrong result";
}
#endif
run_is_sorted_until_all_scenarios<DynamicTag, double>();
run_is_sorted_until_all_scenarios<StridedTwoTag, double>();
run_is_sorted_until_all_scenarios<StridedThreeTag, double>();

26
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsLexicographicalCompare.cpp
View File

@ -44,16 +44,16 @@ void test_lexicographical_compare(const ViewType1 view_1, ViewType2 view_2) {
std::lexicographical_compare(h_first_1, h_last_1, h_first_2, h_last_2);
// pass iterators
EXPECT_EQ(std_result, KE::lexicographical_compare(exespace(), first_1,
ASSERT_EQ(std_result, KE::lexicographical_compare(exespace(), first_1,
last_1, first_2, last_2));
EXPECT_EQ(std_result,
ASSERT_EQ(std_result,
KE::lexicographical_compare("label", exespace(), first_1, last_1,
first_2, last_2));
// pass views
EXPECT_EQ(std_result,
ASSERT_EQ(std_result,
KE::lexicographical_compare(exespace(), view_1, view_2));
EXPECT_EQ(std_result,
ASSERT_EQ(std_result,
KE::lexicographical_compare("label", exespace(), view_1, view_2));
}
@ -67,17 +67,17 @@ void test_lexicographical_compare(const ViewType1 view_1, ViewType2 view_2) {
h_first_1, h_last_1, h_first_2, h_last_2, custom_comparator);
// pass iterators
EXPECT_EQ(std_result,
ASSERT_EQ(std_result,
KE::lexicographical_compare(exespace(), first_1, last_1, first_2,
last_2, custom_comparator));
EXPECT_EQ(std_result,
ASSERT_EQ(std_result,
KE::lexicographical_compare("label", exespace(), first_1, last_1,
first_2, last_2, custom_comparator));
// pass views
EXPECT_EQ(std_result, KE::lexicographical_compare(
ASSERT_EQ(std_result, KE::lexicographical_compare(
exespace(), view_1, view_2, custom_comparator));
EXPECT_EQ(std_result,
ASSERT_EQ(std_result,
KE::lexicographical_compare("label", exespace(), view_1, view_2,
custom_comparator));
}
@ -86,7 +86,7 @@ void test_lexicographical_compare(const ViewType1 view_1, ViewType2 view_2) {
// empty vs non-empty
auto std_result =
std::lexicographical_compare(h_first_1, h_first_1, h_first_2, h_last_2);
EXPECT_EQ(std_result, KE::lexicographical_compare(
ASSERT_EQ(std_result, KE::lexicographical_compare(
exespace(), first_1, first_1, first_2, last_2));
}
@ -95,7 +95,7 @@ void test_lexicographical_compare(const ViewType1 view_1, ViewType2 view_2) {
if (view_1.extent(0) > 1) {
auto std_result = std::lexicographical_compare(h_first_1, h_last_1 - 1,
h_first_2, h_last_2);
EXPECT_EQ(std_result,
ASSERT_EQ(std_result,
KE::lexicographical_compare(exespace(), first_1, last_1 - 1,
first_2, last_2));
}
@ -140,12 +140,6 @@ void run_all_scenarios() {
}
TEST(std_algorithms_lexicographical_compare_test, test) {
#if defined(KOKKOS_ENABLE_CUDA) && \
defined(KOKKOS_COMPILER_NVHPC) // FIXME_NVHPC
if constexpr (std::is_same_v<exespace, Kokkos::Cuda>) {
GTEST_SKIP() << "FIXME wrong result";
}
#endif
// FIXME: should this disable only custom comparator tests?
#if !defined KOKKOS_ENABLE_OPENMPTARGET
run_all_scenarios<DynamicTag, double>();

20
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsMinMaxElementOps.cpp
View File

@ -173,7 +173,7 @@ void std_algo_min_max_test_verify(Kokkos::pair<IndexType, ValueType> goldPair,
const ItType result,
TestedViewType testedView) {
// check that iterator is pointing to right element
EXPECT_EQ(result - KE::begin(testedView), goldPair.first);
ASSERT_EQ(result - KE::begin(testedView), goldPair.first);
// create a view for the result to copy into it the iterator's value
using result_view_t = Kokkos::View<int>;
@ -184,7 +184,7 @@ void std_algo_min_max_test_verify(Kokkos::pair<IndexType, ValueType> goldPair,
Kokkos::create_mirror_view_and_copy(Kokkos::HostSpace(), resultView);
// use the host mirror of the result view to check that the values match
EXPECT_EQ(result_v_h(), goldPair.second);
ASSERT_EQ(result_v_h(), goldPair.second);
}
template <class GoldSolutionType, class ItType, class TestedViewType>
@ -199,39 +199,39 @@ template <class ViewType>
void test_max_element_trivial_data(ViewType view) {
/* if we pass empty range, should return last */
auto result = KE::max_element(exespace(), KE::cbegin(view), KE::cbegin(view));
EXPECT_EQ(result, KE::cbegin(view));
ASSERT_EQ(result, KE::cbegin(view));
/* if we pass empty range, should return last */
auto it0 = KE::cbegin(view) + 3;
auto it1 = it0;
auto result2 = KE::max_element(exespace(), it0, it1);
EXPECT_EQ(result2, it1);
ASSERT_EQ(result2, it1);
}
template <class ViewType>
void test_min_element_trivial_data(ViewType view) {
/* if we pass empty range, should return last */
auto result = KE::min_element(exespace(), KE::cbegin(view), KE::cbegin(view));
EXPECT_EQ(result, KE::cbegin(view));
ASSERT_EQ(result, KE::cbegin(view));
/* if we pass empty range, should return last */
auto it0 = KE::cbegin(view) + 3;
auto it1 = it0;
auto result2 = KE::min_element(exespace(), it0, it1);
EXPECT_EQ(result2, it1);
ASSERT_EQ(result2, it1);
}
template <class ViewType>
void test_minmax_element_empty_range(ViewType view) {
auto result =
KE::minmax_element(exespace(), KE::cbegin(view), KE::cbegin(view));
EXPECT_EQ(result.first, KE::cbegin(view));
EXPECT_EQ(result.second, KE::cbegin(view));
ASSERT_EQ(result.first, KE::cbegin(view));
ASSERT_EQ(result.second, KE::cbegin(view));
auto it0 = KE::cbegin(view) + 3;
auto it1 = it0;
auto result2 = KE::minmax_element(exespace(), it0, it1);
EXPECT_EQ(result2.first, it1);
EXPECT_EQ(result2.second, it1);
ASSERT_EQ(result2.first, it1);
ASSERT_EQ(result2.second, it1);
}
template <class ViewType>

22
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsMismatch.cpp
View File

@ -120,10 +120,10 @@ void run_single_scenario(ViewType view1, ViewType view2,
const auto my_diff12 = my_res1.second - f2;
const auto my_diff21 = my_res2.first - f1;
const auto my_diff22 = my_res2.second - f2;
EXPECT_EQ(my_diff11, std_diff1);
EXPECT_EQ(my_diff12, std_diff2);
EXPECT_EQ(my_diff21, std_diff1);
EXPECT_EQ(my_diff22, std_diff2);
ASSERT_EQ(my_diff11, std_diff1);
ASSERT_EQ(my_diff12, std_diff2);
ASSERT_EQ(my_diff21, std_diff1);
ASSERT_EQ(my_diff22, std_diff2);
}
{
@ -134,10 +134,10 @@ void run_single_scenario(ViewType view1, ViewType view2,
const auto my_diff12 = my_res1.second - KE::begin(view2);
const auto my_diff21 = my_res2.first - KE::begin(view1);
const auto my_diff22 = my_res2.second - KE::begin(view2);
EXPECT_EQ(my_diff11, std_diff1);
EXPECT_EQ(my_diff12, std_diff2);
EXPECT_EQ(my_diff21, std_diff1);
EXPECT_EQ(my_diff22, std_diff2);
ASSERT_EQ(my_diff11, std_diff1);
ASSERT_EQ(my_diff12, std_diff2);
ASSERT_EQ(my_diff21, std_diff1);
ASSERT_EQ(my_diff22, std_diff2);
}
}
@ -189,12 +189,6 @@ void run_all_scenarios() {
}
TEST(std_algorithms_mismatch_test, test) {
#if defined(KOKKOS_ENABLE_CUDA) && \
defined(KOKKOS_COMPILER_NVHPC) // FIXME_NVHPC
if constexpr (std::is_same_v<exespace, Kokkos::Cuda>) {
GTEST_SKIP() << "FIXME wrong result";
}
#endif
run_all_scenarios<DynamicTag, double>();
run_all_scenarios<StridedThreeTag, int>();
}

34
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsModOps.cpp
View File

@ -52,14 +52,14 @@ TEST(std_algorithms_mod_ops_test, move) {
// move constr
MyMovableType b(std::move(a));
EXPECT_EQ(b.m_value, 11);
EXPECT_EQ(a.m_value, -2);
ASSERT_EQ(b.m_value, 11);
ASSERT_EQ(a.m_value, -2);
// move assign
MyMovableType c;
c = std::move(b);
EXPECT_EQ(c.m_value, 11);
EXPECT_EQ(b.m_value, -4);
ASSERT_EQ(c.m_value, 11);
ASSERT_EQ(b.m_value, -4);
}
template <class ViewType>
@ -97,8 +97,8 @@ TEST(std_algorithms_mod_ops_test, swap) {
int a = 1;
int b = 2;
KE::swap(a, b);
EXPECT_EQ(a, 2);
EXPECT_EQ(b, 1);
ASSERT_EQ(a, 2);
ASSERT_EQ(b, 1);
}
{
@ -151,17 +151,17 @@ void test_iter_swap(ViewType view) {
using value_type = typename ViewType::value_type;
auto a_dc = create_deep_copyable_compatible_clone(view);
auto a_h = create_mirror_view_and_copy(Kokkos::HostSpace(), a_dc);
EXPECT_EQ(view.extent_int(0), 10);
EXPECT_EQ(a_h(0), value_type(3));
EXPECT_EQ(a_h(1), value_type(1));
EXPECT_EQ(a_h(2), value_type(2));
EXPECT_EQ(a_h(3), value_type(0));
EXPECT_EQ(a_h(4), value_type(6));
EXPECT_EQ(a_h(5), value_type(5));
EXPECT_EQ(a_h(6), value_type(4));
EXPECT_EQ(a_h(7), value_type(7));
EXPECT_EQ(a_h(8), value_type(8));
EXPECT_EQ(a_h(9), value_type(9));
ASSERT_EQ(view.extent_int(0), 10);
ASSERT_EQ(a_h(0), value_type(3));
ASSERT_EQ(a_h(1), value_type(1));
ASSERT_EQ(a_h(2), value_type(2));
ASSERT_EQ(a_h(3), value_type(0));
ASSERT_EQ(a_h(4), value_type(6));
ASSERT_EQ(a_h(5), value_type(5));
ASSERT_EQ(a_h(6), value_type(4));
ASSERT_EQ(a_h(7), value_type(7));
ASSERT_EQ(a_h(8), value_type(8));
ASSERT_EQ(a_h(9), value_type(9));
}
TEST(std_algorithms_mod_ops_test, iter_swap_static_view) {

86
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsModSeqOps.cpp
View File

@ -34,21 +34,21 @@ struct std_algorithms_mod_seq_ops_test : std_algorithms_test {
TEST_F(std_algorithms_mod_seq_ops_test, copy) {
auto result = KE::copy(exespace(), KE::begin(m_static_view),
KE::end(m_static_view), KE::begin(m_strided_view));
EXPECT_EQ(KE::end(m_strided_view), result);
ASSERT_EQ(KE::end(m_strided_view), result);
compare_views(m_static_view, m_strided_view);
auto result2 = KE::copy(exespace(), KE::begin(m_strided_view),
KE::end(m_strided_view), KE::begin(m_dynamic_view));
EXPECT_EQ(KE::end(m_dynamic_view), result2);
ASSERT_EQ(KE::end(m_dynamic_view), result2);
compare_views(m_dynamic_view, m_strided_view);
}
TEST_F(std_algorithms_mod_seq_ops_test, copy_view) {
EXPECT_EQ(KE::end(m_dynamic_view),
ASSERT_EQ(KE::end(m_dynamic_view),
KE::copy(exespace(), m_static_view, m_dynamic_view));
compare_views(m_static_view, m_dynamic_view);
EXPECT_EQ(KE::end(m_strided_view),
ASSERT_EQ(KE::end(m_strided_view),
KE::copy(exespace(), m_dynamic_view, m_strided_view));
compare_views(m_dynamic_view, m_strided_view);
}
@ -70,11 +70,11 @@ TEST_F(std_algorithms_mod_seq_ops_test, copy_n) {
// pass iterators
auto first = KE::begin(m_static_view);
auto dest = KE::begin(m_dynamic_view);
EXPECT_EQ(dest + n, KE::copy_n(exespace(), first, n, dest));
ASSERT_EQ(dest + n, KE::copy_n(exespace(), first, n, dest));
compare_views(expected, m_dynamic_view);
// pass views
EXPECT_EQ(KE::begin(m_strided_view) + n,
ASSERT_EQ(KE::begin(m_strided_view) + n,
KE::copy_n(exespace(), m_static_view, n, m_strided_view));
compare_views(expected, m_strided_view);
}
@ -85,12 +85,12 @@ TEST_F(std_algorithms_mod_seq_ops_test, copy_backward) {
auto dest = KE::end(m_dynamic_view);
// pass iterators
EXPECT_EQ(KE::begin(m_dynamic_view),
ASSERT_EQ(KE::begin(m_dynamic_view),
KE::copy_backward(exespace(), first, last, dest));
compare_views(m_static_view, m_dynamic_view);
// pass views
EXPECT_EQ(KE::begin(m_strided_view),
ASSERT_EQ(KE::begin(m_strided_view),
KE::copy_backward(exespace(), m_static_view, m_strided_view));
compare_views(m_static_view, m_strided_view);
}
@ -112,11 +112,11 @@ TEST_F(std_algorithms_mod_seq_ops_test, reverse_copy) {
auto last = KE::end(m_static_view);
auto dest = KE::begin(m_dynamic_view);
EXPECT_EQ(KE::end(m_dynamic_view),
ASSERT_EQ(KE::end(m_dynamic_view),
KE::reverse_copy(exespace(), first, last, dest));
compare_views(expected, m_dynamic_view);
EXPECT_EQ(KE::end(m_strided_view),
ASSERT_EQ(KE::end(m_strided_view),
KE::reverse_copy(exespace(), m_static_view, m_strided_view));
compare_views(expected, m_strided_view);
}
@ -151,25 +151,25 @@ TEST_F(std_algorithms_mod_seq_ops_test, fill_n) {
// fill all elements
// pass iterator
EXPECT_EQ(KE::end(m_static_view),
ASSERT_EQ(KE::end(m_static_view),
KE::fill_n(exespace(), KE::begin(m_static_view),
m_static_view.extent(0), fill_n_value));
verify_values(fill_n_value, m_static_view);
// pass view
EXPECT_EQ(KE::end(m_strided_view),
ASSERT_EQ(KE::end(m_strided_view),
KE::fill_n(exespace(), m_strided_view, m_strided_view.extent(0),
fill_n_value));
verify_values(fill_n_value, m_strided_view);
// fill zero elements
// pass view
EXPECT_EQ(KE::begin(m_dynamic_view),
ASSERT_EQ(KE::begin(m_dynamic_view),
KE::fill_n(exespace(), m_dynamic_view, 0, fill_n_new_value));
// fill single element
// pass iterator
EXPECT_EQ(
ASSERT_EQ(
KE::begin(m_static_view) + 1,
KE::fill_n(exespace(), KE::begin(m_static_view), 1, fill_n_new_value));
@ -212,21 +212,21 @@ TEST_F(std_algorithms_mod_seq_ops_test, transform_from_fixture_unary_op) {
auto r1 = KE::transform(exespace(), KE::begin(m_static_view),
KE::end(m_static_view), KE::begin(m_dynamic_view),
TransformFunctor());
EXPECT_EQ(r1, KE::end(m_dynamic_view));
ASSERT_EQ(r1, KE::end(m_dynamic_view));
compare_views(gold_source, m_static_view);
verify_values(-1., m_dynamic_view);
// transform dynamic view, store results in strided view
auto r2 = KE::transform(exespace(), m_dynamic_view, m_strided_view,
TransformFunctor());
EXPECT_EQ(r2, KE::end(m_strided_view));
ASSERT_EQ(r2, KE::end(m_strided_view));
verify_values(-1., m_dynamic_view);
verify_values(-1., m_strided_view);
// transform strided view, store results in static view
auto r3 = KE::transform(exespace(), m_strided_view, m_static_view,
TransformFunctor());
EXPECT_EQ(r3, KE::end(m_static_view));
ASSERT_EQ(r3, KE::end(m_static_view));
verify_values(-1., m_static_view);
verify_values(-1., m_strided_view);
}
@ -254,7 +254,7 @@ TEST_F(std_algorithms_mod_seq_ops_test, transform_from_fixture_binary_op) {
auto r1 = KE::transform(exespace(), KE::begin(m_static_view),
KE::end(m_static_view), KE::begin(m_dynamic_view),
KE::begin(m_strided_view), TransformBinaryFunctor());
EXPECT_EQ(r1, KE::end(m_strided_view));
ASSERT_EQ(r1, KE::end(m_strided_view));
compare_views(expected, m_strided_view);
expected(0) = 0;
@ -269,7 +269,7 @@ TEST_F(std_algorithms_mod_seq_ops_test, transform_from_fixture_binary_op) {
expected(9) = 18;
auto r2 = KE::transform("label", exespace(), m_static_view, m_strided_view,
m_dynamic_view, TransformBinaryFunctor());
EXPECT_EQ(r2, KE::end(m_dynamic_view));
ASSERT_EQ(r2, KE::end(m_dynamic_view));
compare_views(expected, m_dynamic_view);
}
@ -296,19 +296,19 @@ TEST_F(std_algorithms_mod_seq_ops_test, generate) {
TEST_F(std_algorithms_mod_seq_ops_test, generate_n) {
// iterator + functor
EXPECT_EQ(KE::end(m_static_view),
ASSERT_EQ(KE::end(m_static_view),
KE::generate_n(exespace(), KE::begin(m_static_view),
m_static_view.extent(0), GenerateFunctor()));
verify_values(generated_value, m_static_view);
// view + functor
EXPECT_EQ(KE::end(m_dynamic_view),
ASSERT_EQ(KE::end(m_dynamic_view),
KE::generate_n(exespace(), m_dynamic_view, m_dynamic_view.extent(0),
GenerateFunctor()));
verify_values(generated_value, m_dynamic_view);
// view + functor, negative n
EXPECT_EQ(KE::begin(m_strided_view),
ASSERT_EQ(KE::begin(m_strided_view),
KE::generate_n(exespace(), m_strided_view, -1, GenerateFunctor()));
}
@ -352,7 +352,7 @@ void test_swap_ranges(ViewType view) {
auto last1 = first1 + 4;
auto first2 = KE::begin(viewB) + 1;
auto r = KE::swap_ranges(exespace(), first1, last1, first2);
EXPECT_EQ(r, first2 + 4);
ASSERT_EQ(r, first2 + 4);
/* check VIEW_A */
static_view_type checkViewA("tmp");
@ -360,16 +360,16 @@ void test_swap_ranges(ViewType view) {
parallel_for(ext, cp_func_a_t(view, checkViewA));
auto cvA_h =
Kokkos::create_mirror_view_and_copy(Kokkos::HostSpace(), checkViewA);
EXPECT_EQ(cvA_h(0), 0);
EXPECT_EQ(cvA_h(1), 1);
EXPECT_EQ(cvA_h(2), 99);
EXPECT_EQ(cvA_h(3), 98);
EXPECT_EQ(cvA_h(4), 97);
EXPECT_EQ(cvA_h(5), 96);
EXPECT_EQ(cvA_h(6), 6);
EXPECT_EQ(cvA_h(7), 7);
EXPECT_EQ(cvA_h(8), 8);
EXPECT_EQ(cvA_h(9), 9);
ASSERT_EQ(cvA_h(0), 0);
ASSERT_EQ(cvA_h(1), 1);
ASSERT_EQ(cvA_h(2), 99);
ASSERT_EQ(cvA_h(3), 98);
ASSERT_EQ(cvA_h(4), 97);
ASSERT_EQ(cvA_h(5), 96);
ASSERT_EQ(cvA_h(6), 6);
ASSERT_EQ(cvA_h(7), 7);
ASSERT_EQ(cvA_h(8), 8);
ASSERT_EQ(cvA_h(9), 9);
/* check viewB */
static_view_type checkViewB("tmpB");
@ -377,16 +377,16 @@ void test_swap_ranges(ViewType view) {
Kokkos::parallel_for(ext, cp_func_b_t(viewB, checkViewB));
auto cvB_h =
Kokkos::create_mirror_view_and_copy(Kokkos::HostSpace(), checkViewB);
EXPECT_EQ(cvB_h(0), 100);
EXPECT_EQ(cvB_h(1), 2);
EXPECT_EQ(cvB_h(2), 3);
EXPECT_EQ(cvB_h(3), 4);
EXPECT_EQ(cvB_h(4), 5);
EXPECT_EQ(cvB_h(5), 95);
EXPECT_EQ(cvB_h(6), 94);
EXPECT_EQ(cvB_h(7), 93);
EXPECT_EQ(cvB_h(8), 92);
EXPECT_EQ(cvB_h(9), 91);
ASSERT_EQ(cvB_h(0), 100);
ASSERT_EQ(cvB_h(1), 2);
ASSERT_EQ(cvB_h(2), 3);
ASSERT_EQ(cvB_h(3), 4);
ASSERT_EQ(cvB_h(4), 5);
ASSERT_EQ(cvB_h(5), 95);
ASSERT_EQ(cvB_h(6), 94);
ASSERT_EQ(cvB_h(7), 93);
ASSERT_EQ(cvB_h(8), 92);
ASSERT_EQ(cvB_h(9), 91);
}
TEST_F(std_algorithms_mod_seq_ops_test, swap_ranges) {

8
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsMoveBackward.cpp
View File

@ -53,20 +53,20 @@ void run_single_scenario(const InfoType& scenario_info, int apiId) {
auto rit =
KE::move_backward(exespace(), KE::begin(v), KE::end(v), KE::end(v2));
const int dist = KE::distance(KE::begin(v2), rit);
EXPECT_EQ(dist, 5);
ASSERT_EQ(dist, 5);
} else if (apiId == 1) {
auto rit = KE::move_backward("mylabel", exespace(), KE::begin(v),
KE::end(v), KE::end(v2));
const int dist = KE::distance(KE::begin(v2), rit);
EXPECT_EQ(dist, 5);
ASSERT_EQ(dist, 5);
} else if (apiId == 2) {
auto rit = KE::move_backward(exespace(), v, v2);
const int dist = KE::distance(KE::begin(v2), rit);
EXPECT_EQ(dist, 5);
ASSERT_EQ(dist, 5);
} else if (apiId == 3) {
auto rit = KE::move_backward("mylabel", exespace(), v, v2);
const int dist = KE::distance(KE::begin(v2), rit);
EXPECT_EQ(dist, 5);
ASSERT_EQ(dist, 5);
}
// check

72
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsNumerics.cpp
View File

@ -151,8 +151,8 @@ void run_and_check_transform_reduce_default(ViewType1 first_view,
const auto r2 = KE::transform_reduce(
"MYLABEL", ExecutionSpace(), KE::cbegin(first_view),
KE::cbegin(first_view), KE::cbegin(second_view), init_value);
EXPECT_EQ(r1, init_value);
EXPECT_EQ(r2, init_value);
ASSERT_EQ(r1, init_value);
ASSERT_EQ(r2, init_value);
// non-trivial cases
const auto r3 = KE::transform_reduce(ExecutionSpace(), KE::cbegin(first_view),
@ -168,10 +168,10 @@ void run_and_check_transform_reduce_default(ViewType1 first_view,
const auto r6 = KE::transform_reduce("MYLABEL", ExecutionSpace(), first_view,
second_view, init_value);
EXPECT_EQ(r3, result_value);
EXPECT_EQ(r4, result_value);
EXPECT_EQ(r5, result_value);
EXPECT_EQ(r6, result_value);
ASSERT_EQ(r3, result_value);
ASSERT_EQ(r4, result_value);
ASSERT_EQ(r5, result_value);
ASSERT_EQ(r6, result_value);
}
TEST_F(std_algorithms_numerics_test,
@ -254,8 +254,8 @@ void run_and_check_transform_reduce_overloadA(ViewType1 first_view,
KE::cbegin(first_view), KE::cbegin(second_view),
init_value, std::forward<Args>(args)...);
EXPECT_EQ(r1, init_value);
EXPECT_EQ(r2, init_value);
ASSERT_EQ(r1, init_value);
ASSERT_EQ(r2, init_value);
// non trivial cases
const auto r3 = KE::transform_reduce(
@ -273,10 +273,10 @@ void run_and_check_transform_reduce_overloadA(ViewType1 first_view,
KE::transform_reduce("MYLABEL", ExecutionSpace(), first_view, second_view,
init_value, std::forward<Args>(args)...);
EXPECT_EQ(r3, result_value);
EXPECT_EQ(r4, result_value);
EXPECT_EQ(r5, result_value);
EXPECT_EQ(r6, result_value);
ASSERT_EQ(r3, result_value);
ASSERT_EQ(r4, result_value);
ASSERT_EQ(r5, result_value);
ASSERT_EQ(r6, result_value);
}
TEST_F(std_algorithms_numerics_test,
@ -373,8 +373,8 @@ void run_and_check_transform_reduce_overloadB(ViewType view,
KE::cbegin(view), KE::cbegin(view),
init_value, std::forward<Args>(args)...);
EXPECT_EQ(r1, init_value);
EXPECT_EQ(r2, init_value);
ASSERT_EQ(r1, init_value);
ASSERT_EQ(r2, init_value);
// non trivial
const auto r3 =
@ -390,10 +390,10 @@ void run_and_check_transform_reduce_overloadB(ViewType view,
const auto r6 = KE::transform_reduce("MYLABEL", ExecutionSpace(), view,
init_value, std::forward<Args>(args)...);
EXPECT_EQ(r3, result_value);
EXPECT_EQ(r4, result_value);
EXPECT_EQ(r5, result_value);
EXPECT_EQ(r6, result_value);
ASSERT_EQ(r3, result_value);
ASSERT_EQ(r4, result_value);
ASSERT_EQ(r5, result_value);
ASSERT_EQ(r6, result_value);
}
TEST_F(std_algorithms_numerics_test,
@ -447,8 +447,8 @@ void run_and_check_reduce_overloadA(ViewType view, ValueType non_trivial_result,
KE::reduce(ExecutionSpace(), KE::cbegin(view), KE::cbegin(view));
const auto r2 = KE::reduce("MYLABEL", ExecutionSpace(), KE::cbegin(view),
KE::cbegin(view));
EXPECT_EQ(r1, trivial_result);
EXPECT_EQ(r2, trivial_result);
ASSERT_EQ(r1, trivial_result);
ASSERT_EQ(r2, trivial_result);
// non trivial cases
const auto r3 =
@ -458,10 +458,10 @@ void run_and_check_reduce_overloadA(ViewType view, ValueType non_trivial_result,
const auto r5 = KE::reduce(ExecutionSpace(), view);
const auto r6 = KE::reduce("MYLABEL", ExecutionSpace(), view);
EXPECT_EQ(r3, non_trivial_result);
EXPECT_EQ(r4, non_trivial_result);
EXPECT_EQ(r5, non_trivial_result);
EXPECT_EQ(r6, non_trivial_result);
ASSERT_EQ(r3, non_trivial_result);
ASSERT_EQ(r4, non_trivial_result);
ASSERT_EQ(r5, non_trivial_result);
ASSERT_EQ(r6, non_trivial_result);
}
TEST_F(std_algorithms_numerics_test,
@ -503,8 +503,8 @@ void run_and_check_reduce_overloadB(ViewType view, ValueType result_value,
KE::cbegin(view), init_value);
const auto r2 = KE::reduce("MYLABEL", ExecutionSpace(), KE::cbegin(view),
KE::cbegin(view), init_value);
EXPECT_EQ(r1, init_value);
EXPECT_EQ(r2, init_value);
ASSERT_EQ(r1, init_value);
ASSERT_EQ(r2, init_value);
// non trivial cases
const auto r3 = KE::reduce(ExecutionSpace(), KE::cbegin(view), KE::cend(view),
@ -514,10 +514,10 @@ void run_and_check_reduce_overloadB(ViewType view, ValueType result_value,
const auto r5 = KE::reduce(ExecutionSpace(), view, init_value);
const auto r6 = KE::reduce("MYLABEL", ExecutionSpace(), view, init_value);
EXPECT_EQ(r3, result_value);
EXPECT_EQ(r4, result_value);
EXPECT_EQ(r5, result_value);
EXPECT_EQ(r6, result_value);
ASSERT_EQ(r3, result_value);
ASSERT_EQ(r4, result_value);
ASSERT_EQ(r5, result_value);
ASSERT_EQ(r6, result_value);
}
TEST_F(std_algorithms_numerics_test,
@ -553,8 +553,8 @@ void run_and_check_reduce_overloadC(ViewType view, ValueType result_value,
KE::cbegin(view), init_value, joiner);
const auto r2 = KE::reduce("MYLABEL", ExecutionSpace(), KE::cbegin(view),
KE::cbegin(view), init_value, joiner);
EXPECT_EQ(r1, init_value);
EXPECT_EQ(r2, init_value);
ASSERT_EQ(r1, init_value);
ASSERT_EQ(r2, init_value);
// non trivial cases
const auto r3 = KE::reduce(ExecutionSpace(), KE::cbegin(view), KE::cend(view),
@ -565,10 +565,10 @@ void run_and_check_reduce_overloadC(ViewType view, ValueType result_value,
const auto r6 =
KE::reduce("MYLABEL", ExecutionSpace(), view, init_value, joiner);
EXPECT_EQ(r3, result_value);
EXPECT_EQ(r4, result_value);
EXPECT_EQ(r5, result_value);
EXPECT_EQ(r6, result_value);
ASSERT_EQ(r3, result_value);
ASSERT_EQ(r4, result_value);
ASSERT_EQ(r5, result_value);
ASSERT_EQ(r6, result_value);
}
TEST_F(std_algorithms_numerics_test,

36
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsPartitionCopy.cpp
View File

@ -130,12 +130,12 @@ void verify_data(const std::string& name, ResultType my_result,
const std::size_t my_diff_true = my_result.first - KE::begin(view_dest_true);
const std::size_t my_diff_false =
my_result.second - KE::begin(view_dest_false);
EXPECT_EQ(std_diff_true, my_diff_true);
EXPECT_EQ(std_diff_false, my_diff_false);
ASSERT_EQ(std_diff_true, my_diff_true);
ASSERT_EQ(std_diff_false, my_diff_false);
auto view_dest_true_h = create_host_space_copy(view_dest_true);
for (std::size_t i = 0; i < std_diff_true; ++i) {
EXPECT_EQ(std_vec_true[i], view_dest_true_h(i));
ASSERT_EQ(std_vec_true[i], view_dest_true_h(i));
// std::cout << "i= " << i << " "
// << " std_true = " << std_vec_true[i] << " "
// << " mine = " << view_dest_true_h(i) << '\n';
@ -143,45 +143,45 @@ void verify_data(const std::string& name, ResultType my_result,
auto view_dest_false_h = create_host_space_copy(view_dest_false);
for (std::size_t i = 0; i < std_diff_false; ++i) {
EXPECT_EQ(std_vec_false[i], view_dest_false_h(i));
ASSERT_EQ(std_vec_false[i], view_dest_false_h(i));
// std::cout << "i= " << i << " "
// << " std_false = " << std_vec_false[i] << " "
// << " mine = " << view_dest_false_h(i) << '\n';
}
if (name == "empty") {
EXPECT_EQ(my_diff_true, 0u);
EXPECT_EQ(my_diff_false, 0u);
ASSERT_EQ(my_diff_true, 0u);
ASSERT_EQ(my_diff_false, 0u);
}
else if (name == "one-element-a") {
EXPECT_EQ(my_diff_true, 0u);
EXPECT_EQ(my_diff_false, 1u);
ASSERT_EQ(my_diff_true, 0u);
ASSERT_EQ(my_diff_false, 1u);
}
else if (name == "one-element-b") {
EXPECT_EQ(my_diff_true, 1u);
EXPECT_EQ(my_diff_false, 0u);
ASSERT_EQ(my_diff_true, 1u);
ASSERT_EQ(my_diff_false, 0u);
}
else if (name == "two-elements-a") {
EXPECT_EQ(my_diff_true, 1u);
EXPECT_EQ(my_diff_false, 1u);
ASSERT_EQ(my_diff_true, 1u);
ASSERT_EQ(my_diff_false, 1u);
}
else if (name == "two-elements-b") {
EXPECT_EQ(my_diff_true, 1u);
EXPECT_EQ(my_diff_false, 1u);
ASSERT_EQ(my_diff_true, 1u);
ASSERT_EQ(my_diff_false, 1u);
}
else if (name == "small-b") {
EXPECT_EQ(my_diff_true, 13u);
EXPECT_EQ(my_diff_false, 0u);
ASSERT_EQ(my_diff_true, 13u);
ASSERT_EQ(my_diff_false, 0u);
}
else if (name == "small-c") {
EXPECT_EQ(my_diff_true, 0u);
EXPECT_EQ(my_diff_false, 15u);
ASSERT_EQ(my_diff_true, 0u);
ASSERT_EQ(my_diff_false, 15u);
}
}

42
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsPartitioningOps.cpp
View File

@ -148,12 +148,6 @@ struct std_algorithms_partitioning_test : public std_algorithms_test {
};
TEST_F(std_algorithms_partitioning_test, is_partitioned_trivial) {
#if defined(KOKKOS_ENABLE_CUDA) && \
defined(KOKKOS_COMPILER_NVHPC) // FIXME_NVHPC
if constexpr (std::is_same_v<exespace, Kokkos::Cuda>) {
GTEST_SKIP() << "FIXME wrong result";
}
#endif
IsNegativeFunctor<value_type> p;
const auto result1 = KE::is_partitioned(exespace(), KE::cbegin(m_static_view),
KE::cbegin(m_static_view), p);
@ -169,12 +163,6 @@ TEST_F(std_algorithms_partitioning_test, is_partitioned_trivial) {
}
TEST_F(std_algorithms_partitioning_test, is_partitioned_accepting_iterators) {
#if defined(KOKKOS_ENABLE_CUDA) && \
defined(KOKKOS_COMPILER_NVHPC) // FIXME_NVHPC
if constexpr (std::is_same_v<exespace, Kokkos::Cuda>) {
GTEST_SKIP() << "FIXME wrong result";
}
#endif
const IsNegativeFunctor<value_type> p;
for (int id = 0; id < FixtureViews::Count; ++id) {
@ -183,25 +171,19 @@ TEST_F(std_algorithms_partitioning_test, is_partitioned_accepting_iterators) {
goldSolutionIsPartitioned(static_cast<FixtureViews>(id));
const auto result1 = KE::is_partitioned(
exespace(), KE::cbegin(m_static_view), KE::cend(m_static_view), p);
EXPECT_EQ(goldBool, result1);
ASSERT_EQ(goldBool, result1);
const auto result2 = KE::is_partitioned(
exespace(), KE::cbegin(m_dynamic_view), KE::cend(m_dynamic_view), p);
EXPECT_EQ(goldBool, result2);
ASSERT_EQ(goldBool, result2);
const auto result3 = KE::is_partitioned(
exespace(), KE::cbegin(m_strided_view), KE::cend(m_strided_view), p);
EXPECT_EQ(goldBool, result3);
ASSERT_EQ(goldBool, result3);
}
}
TEST_F(std_algorithms_partitioning_test, is_partitioned_accepting_view) {
#if defined(KOKKOS_ENABLE_CUDA) && \
defined(KOKKOS_COMPILER_NVHPC) // FIXME_NVHPC
if constexpr (std::is_same_v<exespace, Kokkos::Cuda>) {
GTEST_SKIP() << "FIXME wrong result";
}
#endif
const IsNegativeFunctor<value_type> p;
for (int id = 0; id < FixtureViews::Count; ++id) {
@ -209,23 +191,17 @@ TEST_F(std_algorithms_partitioning_test, is_partitioned_accepting_view) {
const bool goldBool =
goldSolutionIsPartitioned(static_cast<FixtureViews>(id));
const auto result1 = KE::is_partitioned(exespace(), m_static_view, p);
EXPECT_EQ(goldBool, result1);
ASSERT_EQ(goldBool, result1);
const auto result2 = KE::is_partitioned(exespace(), m_dynamic_view, p);
EXPECT_EQ(goldBool, result2);
ASSERT_EQ(goldBool, result2);
const auto result3 = KE::is_partitioned(exespace(), m_strided_view, p);
EXPECT_EQ(goldBool, result3);
ASSERT_EQ(goldBool, result3);
}
}
TEST_F(std_algorithms_partitioning_test, partition_point) {
#if defined(KOKKOS_ENABLE_CUDA) && \
defined(KOKKOS_COMPILER_NVHPC) // FIXME_NVHPC
if constexpr (std::is_same_v<exespace, Kokkos::Cuda>) {
GTEST_SKIP() << "FIXME wrong result";
}
#endif
const IsNegativeFunctor<value_type> p;
for (int id = 0; id < FixtureViews::Count; ++id) {
@ -235,17 +211,17 @@ TEST_F(std_algorithms_partitioning_test, partition_point) {
auto first1 = KE::cbegin(m_static_view);
auto last1 = KE::cend(m_static_view);
const auto result1 = KE::partition_point(exespace(), first1, last1, p);
EXPECT_EQ(goldIndex, result1 - first1);
ASSERT_EQ(goldIndex, result1 - first1);
auto first2 = KE::cbegin(m_dynamic_view);
auto last2 = KE::cend(m_dynamic_view);
const auto result2 = KE::partition_point(exespace(), first2, last2, p);
EXPECT_EQ(goldIndex, result2 - first2);
ASSERT_EQ(goldIndex, result2 - first2);
auto first3 = KE::cbegin(m_strided_view);
auto last3 = KE::cend(m_strided_view);
const auto result3 = KE::partition_point(exespace(), first3, last3, p);
EXPECT_EQ(goldIndex, result3 - first3);
ASSERT_EQ(goldIndex, result3 - first3);
}
}

4
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsRemove.cpp
View File

@ -117,12 +117,12 @@ void verify_data(ViewTypeData view_data_h, ViewTypeTest view_test,
// check that returned iterators are correct
const std::size_t std_diff = std_result - KE::begin(view_data_h);
const std::size_t my_diff = my_result - KE::begin(view_test);
EXPECT_EQ(std_diff, my_diff);
ASSERT_EQ(std_diff, my_diff);
// check the actual data after algo has been applied
auto view_test_h = create_host_space_copy(view_test);
for (std::size_t i = 0; i < my_diff; ++i) {
EXPECT_EQ(view_test_h(i), view_data_h[i]);
ASSERT_EQ(view_test_h(i), view_data_h[i]);
// std::cout << "i= " << i << " "
// << "mine: " << view_test_h(i) << " "
// << "std: " << view_data_h(i)

4
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsRemoveCopy.cpp
View File

@ -135,12 +135,12 @@ void verify_data(ViewFromType view_from, ViewDestType view_dest,
// check that returned iterators are correct
const std::size_t std_diff = std_result - gold_dest_std.begin();
const std::size_t my_diff = my_result - KE::begin(view_dest);
EXPECT_EQ(std_diff, my_diff);
ASSERT_EQ(std_diff, my_diff);
// check the actual data after algo has been applied
auto view_dest_h = create_host_space_copy(view_dest);
for (std::size_t i = 0; i < my_diff; ++i) {
EXPECT_EQ(view_dest_h(i), gold_dest_std[i]);
ASSERT_EQ(view_dest_h(i), gold_dest_std[i]);
// std::cout << "i= " << i << " "
// << "mine: " << view_dest_h(i) << " "
// << "std: " << gold_dest_std[i]

4
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsRemoveCopyIf.cpp
View File

@ -119,12 +119,12 @@ void verify_data(ViewTypeFrom view_from, ViewTypeDest view_dest,
// check that returned iterators are correct
const std::size_t std_diff = std_result - gold_dest_std.begin();
const std::size_t my_diff = my_result - KE::begin(view_dest);
EXPECT_EQ(std_diff, my_diff);
ASSERT_EQ(std_diff, my_diff);
// check the actual data after algo has been applied
auto view_dest_h = create_host_space_copy(view_dest);
for (std::size_t i = 0; i < my_diff; ++i) {
EXPECT_EQ(view_dest_h(i), gold_dest_std[i]);
ASSERT_EQ(view_dest_h(i), gold_dest_std[i]);
// std::cout << "i= " << i << " "
// << "mine: " << view_dest_h(i) << " "
// << "std: " << gold_dest_std[i]

4
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsRemoveIf.cpp
View File

@ -112,12 +112,12 @@ void verify_data(ViewTypeData view_data_h, ViewTypeTest view_test,
// check that returned iterators are correct
const std::size_t std_diff = std_result - KE::begin(view_data_h);
const std::size_t my_diff = my_result - KE::begin(view_test);
EXPECT_EQ(std_diff, my_diff);
ASSERT_EQ(std_diff, my_diff);
// check the actual data after algo has been applied
auto view_test_h = create_host_space_copy(view_test);
for (std::size_t i = 0; i < my_diff; ++i) {
EXPECT_EQ(view_test_h(i), view_data_h[i]);
ASSERT_EQ(view_test_h(i), view_data_h[i]);
// std::cout << "i= " << i << " "
// << "mine: " << view_test_h(i) << " "
// << "std: " << view_data_h(i)

24
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsReplace.cpp
View File

@ -104,30 +104,30 @@ void verify_data(const std::string& name, ViewType1 test_view,
}
else if (name == "one-element-a") {
EXPECT_EQ(view_h(0), ValueType{1});
ASSERT_EQ(view_h(0), ValueType{1});
}
else if (name == "one-element-b") {
EXPECT_EQ(view_h(0), new_value);
ASSERT_EQ(view_h(0), new_value);
}
else if (name == "two-elements-a") {
EXPECT_EQ(view_h(0), ValueType{1});
EXPECT_EQ(view_h(1), new_value);
ASSERT_EQ(view_h(0), ValueType{1});
ASSERT_EQ(view_h(1), new_value);
}
else if (name == "two-elements-b") {
EXPECT_EQ(view_h(0), new_value);
EXPECT_EQ(view_h(1), ValueType{-1});
ASSERT_EQ(view_h(0), new_value);
ASSERT_EQ(view_h(1), ValueType{-1});
}
else if (name == "small-a") {
for (std::size_t i = 0; i < view_h.extent(0); ++i) {
if (i == 0 || i == 3 || i == 5 || i == 6) {
EXPECT_EQ(view_h(i), new_value);
ASSERT_EQ(view_h(i), new_value);
} else {
const auto gold = ValueType{-5} + static_cast<ValueType>(i + 1);
EXPECT_EQ(view_h(i), gold);
ASSERT_EQ(view_h(i), gold);
}
}
}
@ -135,9 +135,9 @@ void verify_data(const std::string& name, ViewType1 test_view,
else if (name == "small-b") {
for (std::size_t i = 0; i < view_h.extent(0); ++i) {
if (i < 4) {
EXPECT_EQ(view_h(i), ValueType{-1});
ASSERT_EQ(view_h(i), ValueType{-1});
} else {
EXPECT_EQ(view_h(i), new_value);
ASSERT_EQ(view_h(i), new_value);
}
}
}
@ -145,9 +145,9 @@ void verify_data(const std::string& name, ViewType1 test_view,
else if (name == "medium" || name == "large") {
for (std::size_t i = 0; i < view_h.extent(0); ++i) {
if (i % 2 == 0) {
EXPECT_EQ(view_h(i), ValueType{-1});
ASSERT_EQ(view_h(i), ValueType{-1});
} else {
EXPECT_EQ(view_h(i), new_value);
ASSERT_EQ(view_h(i), new_value);
}
}
}

56
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsReplaceCopy.cpp
View File

@ -112,40 +112,40 @@ void verify_data(const std::string& name, ViewTypeFrom view_from,
}
else if (name == "one-element-a") {
EXPECT_EQ(view_from_h(0), ValueType{1});
EXPECT_EQ(view_test_h(0), view_from_h(0));
ASSERT_EQ(view_from_h(0), ValueType{1});
ASSERT_EQ(view_test_h(0), view_from_h(0));
}
else if (name == "one-element-b") {
EXPECT_EQ(view_from_h(0), ValueType{2});
EXPECT_EQ(view_test_h(0), new_value);
ASSERT_EQ(view_from_h(0), ValueType{2});
ASSERT_EQ(view_test_h(0), new_value);
}
else if (name == "two-elements-a") {
EXPECT_EQ(view_from_h(0), ValueType{1});
EXPECT_EQ(view_from_h(1), ValueType{2});
ASSERT_EQ(view_from_h(0), ValueType{1});
ASSERT_EQ(view_from_h(1), ValueType{2});
EXPECT_EQ(view_test_h(0), view_from_h(0));
EXPECT_EQ(view_test_h(1), new_value);
ASSERT_EQ(view_test_h(0), view_from_h(0));
ASSERT_EQ(view_test_h(1), new_value);
}
else if (name == "two-elements-b") {
EXPECT_EQ(view_from_h(0), ValueType{2});
EXPECT_EQ(view_from_h(1), ValueType{-1});
ASSERT_EQ(view_from_h(0), ValueType{2});
ASSERT_EQ(view_from_h(1), ValueType{-1});
EXPECT_EQ(view_test_h(0), new_value);
EXPECT_EQ(view_test_h(1), view_from_h(1));
ASSERT_EQ(view_test_h(0), new_value);
ASSERT_EQ(view_test_h(1), view_from_h(1));
}
else if (name == "small-a") {
for (std::size_t i = 0; i < view_test_h.extent(0); ++i) {
if (i == 0 || i == 3 || i == 5 || i == 6) {
EXPECT_EQ(view_from_h(i), ValueType{2});
EXPECT_EQ(view_test_h(i), new_value);
ASSERT_EQ(view_from_h(i), ValueType{2});
ASSERT_EQ(view_test_h(i), new_value);
} else {
const auto gold = ValueType{-5} + static_cast<ValueType>(i + 1);
EXPECT_EQ(view_from_h(i), gold);
EXPECT_EQ(view_test_h(i), gold);
ASSERT_EQ(view_from_h(i), gold);
ASSERT_EQ(view_test_h(i), gold);
}
}
}
@ -153,11 +153,11 @@ void verify_data(const std::string& name, ViewTypeFrom view_from,
else if (name == "small-b") {
for (std::size_t i = 0; i < view_test_h.extent(0); ++i) {
if (i < 4) {
EXPECT_EQ(view_from_h(i), ValueType{-1});
EXPECT_EQ(view_test_h(i), view_from_h(i));
ASSERT_EQ(view_from_h(i), ValueType{-1});
ASSERT_EQ(view_test_h(i), view_from_h(i));
} else {
EXPECT_EQ(view_from_h(i), ValueType{2});
EXPECT_EQ(view_test_h(i), new_value);
ASSERT_EQ(view_from_h(i), ValueType{2});
ASSERT_EQ(view_test_h(i), new_value);
}
}
}
@ -165,11 +165,11 @@ void verify_data(const std::string& name, ViewTypeFrom view_from,
else if (name == "medium" || name == "large") {
for (std::size_t i = 0; i < view_test_h.extent(0); ++i) {
if (i % 2 == 0) {
EXPECT_EQ(view_from_h(i), ValueType{-1});
EXPECT_EQ(view_test_h(i), view_from_h(i));
ASSERT_EQ(view_from_h(i), ValueType{-1});
ASSERT_EQ(view_test_h(i), view_from_h(i));
} else {
EXPECT_EQ(view_from_h(i), ValueType{2});
EXPECT_EQ(view_test_h(i), new_value);
ASSERT_EQ(view_from_h(i), ValueType{2});
ASSERT_EQ(view_test_h(i), new_value);
}
}
}
@ -202,7 +202,7 @@ void run_single_scenario(const InfoType& scenario_info) {
KE::replace_copy(exespace(), KE::cbegin(view_from), KE::cend(view_from),
KE::begin(view_dest), old_value, new_value);
verify_data(name, view_from, view_dest, new_value);
EXPECT_EQ(rit, (KE::begin(view_dest) + view_ext));
ASSERT_EQ(rit, (KE::begin(view_dest) + view_ext));
}
{
@ -215,7 +215,7 @@ void run_single_scenario(const InfoType& scenario_info) {
KE::cend(view_from), KE::begin(view_dest),
old_value, new_value);
verify_data(name, view_from, view_dest, new_value);
EXPECT_EQ(rit, (KE::begin(view_dest) + view_ext));
ASSERT_EQ(rit, (KE::begin(view_dest) + view_ext));
}
{
@ -227,7 +227,7 @@ void run_single_scenario(const InfoType& scenario_info) {
auto rit = KE::replace_copy(exespace(), view_from, view_dest, old_value,
new_value);
verify_data(name, view_from, view_dest, new_value);
EXPECT_EQ(rit, (KE::begin(view_dest) + view_ext));
ASSERT_EQ(rit, (KE::begin(view_dest) + view_ext));
}
{
@ -239,7 +239,7 @@ void run_single_scenario(const InfoType& scenario_info) {
auto rit = KE::replace_copy("label", exespace(), view_from, view_dest,
old_value, new_value);
verify_data(name, view_from, view_dest, new_value);
EXPECT_EQ(rit, (KE::begin(view_dest) + view_ext));
ASSERT_EQ(rit, (KE::begin(view_dest) + view_ext));
}
Kokkos::fence();

56
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsReplaceCopyIf.cpp
View File

@ -112,40 +112,40 @@ void verify_data(const std::string& name, ViewTypeFrom view_from,
}
else if (name == "one-element-a") {
EXPECT_EQ(view_from_h(0), ValueType{1});
EXPECT_EQ(view_test_h(0), view_from_h(0));
ASSERT_EQ(view_from_h(0), ValueType{1});
ASSERT_EQ(view_test_h(0), view_from_h(0));
}
else if (name == "one-element-b") {
EXPECT_EQ(view_from_h(0), ValueType{2});
EXPECT_EQ(view_test_h(0), new_value);
ASSERT_EQ(view_from_h(0), ValueType{2});
ASSERT_EQ(view_test_h(0), new_value);
}
else if (name == "two-elements-a") {
EXPECT_EQ(view_from_h(0), ValueType{1});
EXPECT_EQ(view_from_h(1), ValueType{2});
ASSERT_EQ(view_from_h(0), ValueType{1});
ASSERT_EQ(view_from_h(1), ValueType{2});
EXPECT_EQ(view_test_h(0), view_from_h(0));
EXPECT_EQ(view_test_h(1), new_value);
ASSERT_EQ(view_test_h(0), view_from_h(0));
ASSERT_EQ(view_test_h(1), new_value);
}
else if (name == "two-elements-b") {
EXPECT_EQ(view_from_h(0), ValueType{2});
EXPECT_EQ(view_from_h(1), ValueType{-1});
ASSERT_EQ(view_from_h(0), ValueType{2});
ASSERT_EQ(view_from_h(1), ValueType{-1});
EXPECT_EQ(view_test_h(0), new_value);
EXPECT_EQ(view_test_h(1), view_from_h(1));
ASSERT_EQ(view_test_h(0), new_value);
ASSERT_EQ(view_test_h(1), view_from_h(1));
}
else if (name == "small-a") {
for (std::size_t i = 0; i < view_test_h.extent(0); ++i) {
if (i == 0 || i == 3 || i == 5 || i == 6) {
EXPECT_EQ(view_from_h(i), ValueType{2});
EXPECT_EQ(view_test_h(i), new_value);
ASSERT_EQ(view_from_h(i), ValueType{2});
ASSERT_EQ(view_test_h(i), new_value);
} else {
const auto gold = ValueType{-5} + static_cast<ValueType>(i + 1);
EXPECT_EQ(view_from_h(i), gold);
EXPECT_EQ(view_test_h(i), gold);
ASSERT_EQ(view_from_h(i), gold);
ASSERT_EQ(view_test_h(i), gold);
}
}
}
@ -153,11 +153,11 @@ void verify_data(const std::string& name, ViewTypeFrom view_from,
else if (name == "small-b") {
for (std::size_t i = 0; i < view_test_h.extent(0); ++i) {
if (i < 4) {
EXPECT_EQ(view_from_h(i), ValueType{-1});
EXPECT_EQ(view_test_h(i), view_from_h(i));
ASSERT_EQ(view_from_h(i), ValueType{-1});
ASSERT_EQ(view_test_h(i), view_from_h(i));
} else {
EXPECT_EQ(view_from_h(i), ValueType{2});
EXPECT_EQ(view_test_h(i), new_value);
ASSERT_EQ(view_from_h(i), ValueType{2});
ASSERT_EQ(view_test_h(i), new_value);
}
}
}
@ -165,11 +165,11 @@ void verify_data(const std::string& name, ViewTypeFrom view_from,
else if (name == "medium" || name == "large") {
for (std::size_t i = 0; i < view_test_h.extent(0); ++i) {
if (i % 2 == 0) {
EXPECT_EQ(view_from_h(i), ValueType{-1});
EXPECT_EQ(view_test_h(i), view_from_h(i));
ASSERT_EQ(view_from_h(i), ValueType{-1});
ASSERT_EQ(view_test_h(i), view_from_h(i));
} else {
EXPECT_EQ(view_from_h(i), ValueType{2});
EXPECT_EQ(view_test_h(i), new_value);
ASSERT_EQ(view_from_h(i), ValueType{2});
ASSERT_EQ(view_test_h(i), new_value);
}
}
}
@ -209,7 +209,7 @@ void run_single_scenario(const InfoType& scenario_info) {
KE::cend(view_from), KE::begin(view_dest),
pred_type(), new_value);
verify_data(name, view_from, view_dest, new_value);
EXPECT_EQ(rit, (KE::begin(view_dest) + view_ext));
ASSERT_EQ(rit, (KE::begin(view_dest) + view_ext));
}
{
@ -220,7 +220,7 @@ void run_single_scenario(const InfoType& scenario_info) {
KE::cend(view_from), KE::begin(view_dest),
pred_type(), new_value);
verify_data(name, view_from, view_dest, new_value);
EXPECT_EQ(rit, (KE::begin(view_dest) + view_ext));
ASSERT_EQ(rit, (KE::begin(view_dest) + view_ext));
}
{
@ -230,7 +230,7 @@ void run_single_scenario(const InfoType& scenario_info) {
auto rit = KE::replace_copy_if(exespace(), view_from, view_dest,
pred_type(), new_value);
verify_data(name, view_from, view_dest, new_value);
EXPECT_EQ(rit, (KE::begin(view_dest) + view_ext));
ASSERT_EQ(rit, (KE::begin(view_dest) + view_ext));
}
{
@ -240,7 +240,7 @@ void run_single_scenario(const InfoType& scenario_info) {
auto rit = KE::replace_copy_if("label", exespace(), view_from, view_dest,
pred_type(), new_value);
verify_data(name, view_from, view_dest, new_value);
EXPECT_EQ(rit, (KE::begin(view_dest) + view_ext));
ASSERT_EQ(rit, (KE::begin(view_dest) + view_ext));
}
Kokkos::fence();

2
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsReplaceIf.cpp
View File

@ -138,7 +138,7 @@ void verify_data(ViewType1 data_view, // contains data
// << data_view_dc(i) << " "
// << data_view_h(i) << " "
// << test_view_h(i) << std::endl;
EXPECT_EQ(data_view_h(i), test_view_h(i));
ASSERT_EQ(data_view_h(i), test_view_h(i));
}
}
}

2
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsReverse.cpp
View File

@ -77,7 +77,7 @@ void verify_data(ViewType1 test_view, ViewType2 orig_view) {
const std::size_t ext = test_view.extent(0);
for (std::size_t i = 0; i < ext; ++i) {
EXPECT_EQ(tv_h(i), ov_h(ext - i - 1));
ASSERT_EQ(tv_h(i), ov_h(ext - i - 1));
}
}

4
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsRotate.cpp
View File

@ -136,13 +136,13 @@ void verify_data(ResultIt result_it, ViewType view, ViewHostType data_view_host,
// make sure results match
const auto my_diff = result_it - KE::begin(view);
const auto std_diff = std_rit - KE::begin(data_view_host);
EXPECT_EQ(my_diff, std_diff);
ASSERT_EQ(my_diff, std_diff);
// check views match
auto view_h = create_host_space_copy(view);
const std::size_t ext = view_h.extent(0);
for (std::size_t i = 0; i < ext; ++i) {
EXPECT_EQ(view_h(i), data_view_host[i]);
ASSERT_EQ(view_h(i), data_view_host[i]);
// std::cout << "i= " << i << " "
// << "mine: " << view_h(i) << " "
// << "std: " << data_view_host(i)

10
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsRotateCopy.cpp
View File

@ -139,7 +139,7 @@ void verify_data(ViewTypeFrom view_from, ViewTypeTest view_test,
std_gold_h.begin());
for (std::size_t i = 0; i < ext; ++i) {
EXPECT_EQ(view_test_h(i), std_gold_h[i]);
ASSERT_EQ(view_test_h(i), std_gold_h[i]);
// std::cout << "i= " << i << " "
// << "from: " << view_from_h(i) << " "
// << "mine: " << view_test_h(i) << " "
@ -177,7 +177,7 @@ void run_single_scenario(const InfoType& scenario_info,
auto rit = KE::rotate_copy(exespace(), KE::cbegin(view_from), n_it,
KE::cend(view_from), KE::begin(view_dest));
verify_data(view_from, view_dest, rotation_point);
EXPECT_EQ(rit, (KE::begin(view_dest) + view_ext));
ASSERT_EQ(rit, (KE::begin(view_dest) + view_ext));
}
{
@ -187,7 +187,7 @@ void run_single_scenario(const InfoType& scenario_info,
auto rit = KE::rotate_copy("label", exespace(), KE::cbegin(view_from), n_it,
KE::cend(view_from), KE::begin(view_dest));
verify_data(view_from, view_dest, rotation_point);
EXPECT_EQ(rit, (KE::begin(view_dest) + view_ext));
ASSERT_EQ(rit, (KE::begin(view_dest) + view_ext));
}
{
@ -196,7 +196,7 @@ void run_single_scenario(const InfoType& scenario_info,
auto rit =
KE::rotate_copy(exespace(), view_from, rotation_point, view_dest);
verify_data(view_from, view_dest, rotation_point);
EXPECT_EQ(rit, (KE::begin(view_dest) + view_ext));
ASSERT_EQ(rit, (KE::begin(view_dest) + view_ext));
}
{
@ -205,7 +205,7 @@ void run_single_scenario(const InfoType& scenario_info,
auto rit = KE::rotate_copy("label", exespace(), view_from, rotation_point,
view_dest);
verify_data(view_from, view_dest, rotation_point);
EXPECT_EQ(rit, (KE::begin(view_dest) + view_ext));
ASSERT_EQ(rit, (KE::begin(view_dest) + view_ext));
}
Kokkos::fence();

14
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsSearch.cpp
View File

@ -259,7 +259,7 @@ void run_single_scenario(const InfoType& scenario_info, std::size_t seq_ext,
KE::cbegin(s_view), KE::cend(s_view), args...);
const auto mydiff = myrit - KE::cbegin(view);
const auto stddiff = stdrit - KE::cbegin(view_h);
EXPECT_EQ(mydiff, stddiff);
ASSERT_EQ(mydiff, stddiff);
}
{
@ -268,21 +268,21 @@ void run_single_scenario(const InfoType& scenario_info, std::size_t seq_ext,
KE::cbegin(s_view), KE::cend(s_view), args...);
const auto mydiff = myrit - KE::cbegin(view);
const auto stddiff = stdrit - KE::cbegin(view_h);
EXPECT_EQ(mydiff, stddiff);
ASSERT_EQ(mydiff, stddiff);
}
{
auto myrit = KE::search(exespace(), view, s_view, args...);
const auto mydiff = myrit - KE::begin(view);
const auto stddiff = stdrit - KE::cbegin(view_h);
EXPECT_EQ(mydiff, stddiff);
ASSERT_EQ(mydiff, stddiff);
}
{
auto myrit = KE::search("label", exespace(), view, s_view, args...);
const auto mydiff = myrit - KE::begin(view);
const auto stddiff = stdrit - KE::cbegin(view_h);
EXPECT_EQ(mydiff, stddiff);
ASSERT_EQ(mydiff, stddiff);
}
Kokkos::fence();
@ -325,12 +325,6 @@ void run_all_scenarios() {
}
TEST(std_algorithms_non_mod_seq_ops, search) {
#if defined(KOKKOS_ENABLE_CUDA) && \
defined(KOKKOS_COMPILER_NVHPC) // FIXME_NVHPC
if constexpr (std::is_same_v<exespace, Kokkos::Cuda>) {
GTEST_SKIP() << "FIXME wrong result";
}
#endif
run_all_scenarios<DynamicTag, int>();
run_all_scenarios<StridedThreeTag, int>();
}

14
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsSearch_n.cpp
View File

@ -203,26 +203,26 @@ void run_single_scenario(const InfoType& scenario_info, std::size_t count,
auto myrit = KE::search_n(exespace(), KE::cbegin(view), KE::cend(view),
count, value, args...);
const auto mydiff = myrit - KE::cbegin(view);
EXPECT_EQ(mydiff, stddiff);
ASSERT_EQ(mydiff, stddiff);
}
{
auto myrit = KE::search_n("label", exespace(), KE::cbegin(view),
KE::cend(view), count, value, args...);
const auto mydiff = myrit - KE::cbegin(view);
EXPECT_EQ(mydiff, stddiff);
ASSERT_EQ(mydiff, stddiff);
}
{
auto myrit = KE::search_n("label", exespace(), view, count, value, args...);
const auto mydiff = myrit - KE::begin(view);
EXPECT_EQ(mydiff, stddiff);
ASSERT_EQ(mydiff, stddiff);
}
{
auto myrit = KE::search_n(exespace(), view, count, value, args...);
const auto mydiff = myrit - KE::begin(view);
EXPECT_EQ(mydiff, stddiff);
ASSERT_EQ(mydiff, stddiff);
}
Kokkos::fence();
@ -297,12 +297,6 @@ void run_all_scenarios() {
}
TEST(std_algorithms_non_mod_seq_ops, search_n) {
#if defined(KOKKOS_ENABLE_CUDA) && \
defined(KOKKOS_COMPILER_NVHPC) // FIXME_NVHPC
if constexpr (std::is_same_v<exespace, Kokkos::Cuda>) {
GTEST_SKIP() << "FIXME wrong result";
}
#endif
run_all_scenarios<DynamicTag, int>();
run_all_scenarios<StridedThreeTag, int>();
}

4
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsShiftLeft.cpp
View File

@ -103,12 +103,12 @@ void verify_data(ResultIt result_it, ViewType view, ViewHostType data_view_host,
// make sure results match
const auto my_diff = result_it - KE::begin(view);
const auto std_diff = std_rit - KE::begin(data_view_host);
EXPECT_EQ(my_diff, std_diff);
ASSERT_EQ(my_diff, std_diff);
// check views match
auto view_h = create_host_space_copy(view);
for (std::size_t i = 0; i < (std::size_t)my_diff; ++i) {
EXPECT_EQ(view_h(i), data_view_host[i]);
ASSERT_EQ(view_h(i), data_view_host[i]);
// std::cout << "i= " << i << " "
// << "mine: " << view_h(i) << " "
// << "std: " << data_view_host(i)

4
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsShiftRight.cpp
View File

@ -101,14 +101,14 @@ void verify_data(ResultIt result_it, ViewType view, ViewHostType data_view_host,
// make sure results match
const auto my_diff = KE::end(view) - result_it;
const auto std_diff = KE::end(data_view_host) - std_rit;
EXPECT_EQ(my_diff, std_diff);
ASSERT_EQ(my_diff, std_diff);
// check views match
auto view_h = create_host_space_copy(view);
auto it1 = KE::cbegin(view_h);
auto it2 = KE::cbegin(data_view_host);
for (std::size_t i = 0; i < (std::size_t)my_diff; ++i) {
EXPECT_EQ(it1[i], it2[i]);
ASSERT_EQ(it1[i], it2[i]);
// std::cout << "i= " << i << " "
// << "mine: " << it1[i] << " "
// << "std: " << it2[i]

63
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsTransformExclusiveScan.cpp
View File

@ -165,7 +165,7 @@ void verify_data(ViewType1 data_view, // contains data
// << std::abs(gold_h(i) - test_view_h(i)) << std::endl;
if (std::is_same<gold_view_value_type, int>::value) {
EXPECT_EQ(gold_h(i), test_view_h(i));
ASSERT_EQ(gold_h(i), test_view_h(i));
} else {
const auto error = std::abs(gold_h(i) - test_view_h(i));
if (error > 1e-10) {
@ -221,7 +221,7 @@ void run_single_scenario(const InfoType& scenario_info, ValueType init_value,
auto r = KE::transform_exclusive_scan(
exespace(), KE::cbegin(view_from), KE::cend(view_from),
KE::begin(view_dest), init_value, bop, uop);
EXPECT_EQ(r, KE::end(view_dest));
ASSERT_EQ(r, KE::end(view_dest));
verify_data(view_from, view_dest, init_value, bop, uop);
}
@ -230,7 +230,7 @@ void run_single_scenario(const InfoType& scenario_info, ValueType init_value,
auto r = KE::transform_exclusive_scan(
"label", exespace(), KE::cbegin(view_from), KE::cend(view_from),
KE::begin(view_dest), init_value, bop, uop);
EXPECT_EQ(r, KE::end(view_dest));
ASSERT_EQ(r, KE::end(view_dest));
verify_data(view_from, view_dest, init_value, bop, uop);
}
@ -238,7 +238,7 @@ void run_single_scenario(const InfoType& scenario_info, ValueType init_value,
fill_zero(view_dest);
auto r = KE::transform_exclusive_scan(exespace(), view_from, view_dest,
init_value, bop, uop);
EXPECT_EQ(r, KE::end(view_dest));
ASSERT_EQ(r, KE::end(view_dest));
verify_data(view_from, view_dest, init_value, bop, uop);
}
@ -246,7 +246,7 @@ void run_single_scenario(const InfoType& scenario_info, ValueType init_value,
fill_zero(view_dest);
auto r = KE::transform_exclusive_scan("label", exespace(), view_from,
view_dest, init_value, bop, uop);
EXPECT_EQ(r, KE::end(view_dest));
ASSERT_EQ(r, KE::end(view_dest));
verify_data(view_from, view_dest, init_value, bop, uop);
}
@ -279,6 +279,59 @@ TEST(std_algorithms_numeric_ops_test, transform_exclusive_scan) {
}
#endif
template <class ValueType>
struct MultiplyFunctor {
KOKKOS_INLINE_FUNCTION
ValueType operator()(const ValueType& a, const ValueType& b) const {
return (a * b);
}
};
TEST(std_algorithms_numeric_ops_test, transform_exclusive_scan_functor) {
int dummy = 0;
using view_type = Kokkos::View<int*, exespace>;
view_type dummy_view("dummy_view", 0);
using unary_op_type =
Kokkos::Experimental::Impl::StdNumericScanIdentityReferenceUnaryFunctor<
int>;
using functor_type =
Kokkos::Experimental::Impl::TransformExclusiveScanFunctor<
exespace, int, int, view_type, view_type, MultiplyFunctor<int>,
unary_op_type>;
functor_type functor(dummy, dummy_view, dummy_view, {}, {});
using value_type = functor_type::value_type;
value_type value1;
functor.init(value1);
ASSERT_EQ(value1.val, 0);
ASSERT_EQ(value1.is_initial, true);
value_type value2;
value2.val = 1;
value2.is_initial = false;
functor.join(value1, value2);
ASSERT_EQ(value1.val, 1);
ASSERT_EQ(value1.is_initial, false);
functor.init(value1);
functor.join(value2, value1);
ASSERT_EQ(value2.val, 1);
ASSERT_EQ(value2.is_initial, false);
functor.init(value2);
functor.join(value2, value1);
ASSERT_EQ(value2.val, 0);
ASSERT_EQ(value2.is_initial, true);
value1.val = 3;
value1.is_initial = false;
value2.val = 2;
value2.is_initial = false;
functor.join(value2, value1);
ASSERT_EQ(value2.val, 6);
ASSERT_EQ(value2.is_initial, false);
}
} // namespace TransformEScan
} // namespace stdalgos
} // namespace Test

77
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsTransformInclusiveScan.cpp
View File

@ -177,7 +177,7 @@ void verify_data(ViewType1 data_view, // contains data
// << std::abs(gold_h(i) - test_view_h(i)) << std::endl;
if (std::is_same<gold_view_value_type, int>::value) {
EXPECT_EQ(gold_h(i), test_view_h(i));
ASSERT_EQ(gold_h(i), test_view_h(i));
} else {
const auto error = std::abs(gold_h(i) - test_view_h(i));
if (error > 1e-10) {
@ -246,7 +246,7 @@ void run_single_scenario(const InfoType& scenario_info,
auto r = KE::transform_inclusive_scan(exespace(), KE::cbegin(view_from),
KE::cend(view_from),
KE::begin(view_dest), args...);
EXPECT_EQ(r, KE::end(view_dest));
ASSERT_EQ(r, KE::end(view_dest));
verify_data(view_from, view_dest, args...);
}
@ -255,7 +255,7 @@ void run_single_scenario(const InfoType& scenario_info,
auto r = KE::transform_inclusive_scan(
"label", exespace(), KE::cbegin(view_from), KE::cend(view_from),
KE::begin(view_dest), args...);
EXPECT_EQ(r, KE::end(view_dest));
ASSERT_EQ(r, KE::end(view_dest));
verify_data(view_from, view_dest, args...);
}
@ -263,7 +263,7 @@ void run_single_scenario(const InfoType& scenario_info,
fill_zero(view_dest);
auto r =
KE::transform_inclusive_scan(exespace(), view_from, view_dest, args...);
EXPECT_EQ(r, KE::end(view_dest));
ASSERT_EQ(r, KE::end(view_dest));
verify_data(view_from, view_dest, args...);
}
@ -271,7 +271,7 @@ void run_single_scenario(const InfoType& scenario_info,
fill_zero(view_dest);
auto r = KE::transform_inclusive_scan("label", exespace(), view_from,
view_dest, args...);
EXPECT_EQ(r, KE::end(view_dest));
ASSERT_EQ(r, KE::end(view_dest));
verify_data(view_from, view_dest, args...);
}
@ -306,6 +306,73 @@ TEST(std_algorithms_numeric_ops_test, transform_inclusive_scan) {
}
#endif
template <class ValueType>
struct MultiplyFunctor {
KOKKOS_INLINE_FUNCTION
ValueType operator()(const ValueType& a, const ValueType& b) const {
return (a * b);
}
};
TEST(std_algorithms_numeric_ops_test, transform_inclusive_scan_functor) {
using value_type = KE::Impl::ValueWrapperForNoNeutralElement<int>;
auto test_lambda = [&](auto& functor) {
value_type value1;
functor.init(value1);
ASSERT_EQ(value1.val, 0);
ASSERT_EQ(value1.is_initial, true);
value_type value2;
value2.val = 1;
value2.is_initial = false;
functor.join(value1, value2);
ASSERT_EQ(value1.val, 1);
ASSERT_EQ(value1.is_initial, false);
functor.init(value1);
functor.join(value2, value1);
ASSERT_EQ(value2.val, 1);
ASSERT_EQ(value2.is_initial, false);
functor.init(value2);
functor.join(value2, value1);
ASSERT_EQ(value2.val, 0);
ASSERT_EQ(value2.is_initial, true);
value1.val = 3;
value1.is_initial = false;
value2.val = 2;
value2.is_initial = false;
functor.join(value2, value1);
ASSERT_EQ(value2.val, 6);
ASSERT_EQ(value2.is_initial, false);
};
int dummy = 0;
using view_type = Kokkos::View<int*, exespace>;
view_type dummy_view("dummy_view", 0);
using unary_op_type =
KE::Impl::StdNumericScanIdentityReferenceUnaryFunctor<int>;
{
using functor_type = KE::Impl::TransformInclusiveScanNoInitValueFunctor<
exespace, int, int, view_type, view_type, MultiplyFunctor<int>,
unary_op_type>;
functor_type functor(dummy_view, dummy_view, {}, {});
test_lambda(functor);
}
{
using functor_type = KE::Impl::TransformInclusiveScanWithInitValueFunctor<
exespace, int, int, view_type, view_type, MultiplyFunctor<int>,
unary_op_type>;
functor_type functor(dummy_view, dummy_view, {}, {}, dummy);
test_lambda(functor);
}
}
} // namespace TransformIncScan
} // namespace stdalgos
} // namespace Test

10
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsTransformUnaryOp.cpp
View File

@ -58,7 +58,7 @@ void verify_data(ViewTypeFrom view_from, ViewTypeTest view_test) {
create_mirror_view_and_copy(Kokkos::HostSpace(), view_from_dc);
for (std::size_t i = 0; i < view_test_h.extent(0); ++i) {
EXPECT_EQ(view_test_h(i), view_from_h(i) + value_type(1));
ASSERT_EQ(view_test_h(i), view_from_h(i) + value_type(1));
}
}
@ -89,7 +89,7 @@ void run_single_scenario(const InfoType& scenario_info) {
auto r1 = KE::transform(exespace(), KE::begin(view_from),
KE::end(view_from), KE::begin(view_dest), unOp);
verify_data(view_from, view_dest);
EXPECT_EQ(r1, KE::end(view_dest));
ASSERT_EQ(r1, KE::end(view_dest));
}
{
@ -98,7 +98,7 @@ void run_single_scenario(const InfoType& scenario_info) {
auto r1 = KE::transform("label", exespace(), KE::begin(view_from),
KE::end(view_from), KE::begin(view_dest), unOp);
verify_data(view_from, view_dest);
EXPECT_EQ(r1, KE::end(view_dest));
ASSERT_EQ(r1, KE::end(view_dest));
}
{
@ -106,7 +106,7 @@ void run_single_scenario(const InfoType& scenario_info) {
create_view<ValueType>(Tag{}, view_ext, "transform_uop_dest");
auto r1 = KE::transform(exespace(), view_from, view_dest, unOp);
verify_data(view_from, view_dest);
EXPECT_EQ(r1, KE::end(view_dest));
ASSERT_EQ(r1, KE::end(view_dest));
}
{
@ -114,7 +114,7 @@ void run_single_scenario(const InfoType& scenario_info) {
create_view<ValueType>(Tag{}, view_ext, "transform_uop_dest");
auto r1 = KE::transform("label", exespace(), view_from, view_dest, unOp);
verify_data(view_from, view_dest);
EXPECT_EQ(r1, KE::end(view_dest));
ASSERT_EQ(r1, KE::end(view_dest));
}
Kokkos::fence();

10
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsUnique.cpp
View File

@ -157,7 +157,7 @@ void verify_data(const std::string& name, ResultIt my_result_it,
//
const auto std_diff = (std::size_t)(std_r - KE::begin(data_v_h));
const auto my_diff = (std::size_t)(my_result_it - KE::begin(view_test));
EXPECT_EQ(my_diff, std_diff);
ASSERT_EQ(my_diff, std_diff);
//
// check the data in the view
@ -170,14 +170,14 @@ void verify_data(const std::string& name, ResultIt my_result_it,
// << " my = " << view_test_h(i) << " "
// << " std = " << data_v_h(i)
// << '\n';
EXPECT_EQ(view_test_h(i), data_v_h(i));
ASSERT_EQ(view_test_h(i), data_v_h(i));
}
if (name == "medium-b") {
using value_type = typename ViewType1::value_type;
EXPECT_EQ(my_diff, (std::size_t)2);
EXPECT_EQ(view_test_h(0), (value_type)22);
EXPECT_EQ(view_test_h(1), (value_type)44);
ASSERT_EQ(my_diff, (std::size_t)2);
ASSERT_EQ(view_test_h(0), (value_type)22);
ASSERT_EQ(view_test_h(1), (value_type)44);
}
}

70
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsUniqueCopy.cpp
View File

@ -174,51 +174,51 @@ void verify_data(const std::string& name, ViewTypeFrom view_from,
}
else if (name == "one-element-a") {
EXPECT_EQ(view_test_h(0), static_cast<value_type>(1));
ASSERT_EQ(view_test_h(0), static_cast<value_type>(1));
}
else if (name == "one-element-b") {
EXPECT_EQ(view_test_h(0), static_cast<value_type>(2));
ASSERT_EQ(view_test_h(0), static_cast<value_type>(2));
}
else if (name == "two-elements-a") {
EXPECT_EQ(view_test_h(0), static_cast<value_type>(1));
EXPECT_EQ(view_test_h(1), static_cast<value_type>(2));
ASSERT_EQ(view_test_h(0), static_cast<value_type>(1));
ASSERT_EQ(view_test_h(1), static_cast<value_type>(2));
}
else if (name == "two-elements-b") {
EXPECT_EQ(view_test_h(0), static_cast<value_type>(2));
EXPECT_EQ(view_test_h(1), static_cast<value_type>(-1));
ASSERT_EQ(view_test_h(0), static_cast<value_type>(2));
ASSERT_EQ(view_test_h(1), static_cast<value_type>(-1));
}
else if (name == "small-a") {
EXPECT_EQ(view_test_h(0), static_cast<value_type>(0));
EXPECT_EQ(view_test_h(1), static_cast<value_type>(1));
EXPECT_EQ(view_test_h(2), static_cast<value_type>(2));
EXPECT_EQ(view_test_h(3), static_cast<value_type>(3));
EXPECT_EQ(view_test_h(4), static_cast<value_type>(4));
EXPECT_EQ(view_test_h(5), static_cast<value_type>(5));
EXPECT_EQ(view_test_h(6), static_cast<value_type>(6));
EXPECT_EQ(view_test_h(7), static_cast<value_type>(0));
EXPECT_EQ(view_test_h(8), static_cast<value_type>(0));
EXPECT_EQ(view_test_h(9), static_cast<value_type>(0));
EXPECT_EQ(view_test_h(10), static_cast<value_type>(0));
ASSERT_EQ(view_test_h(0), static_cast<value_type>(0));
ASSERT_EQ(view_test_h(1), static_cast<value_type>(1));
ASSERT_EQ(view_test_h(2), static_cast<value_type>(2));
ASSERT_EQ(view_test_h(3), static_cast<value_type>(3));
ASSERT_EQ(view_test_h(4), static_cast<value_type>(4));
ASSERT_EQ(view_test_h(5), static_cast<value_type>(5));
ASSERT_EQ(view_test_h(6), static_cast<value_type>(6));
ASSERT_EQ(view_test_h(7), static_cast<value_type>(0));
ASSERT_EQ(view_test_h(8), static_cast<value_type>(0));
ASSERT_EQ(view_test_h(9), static_cast<value_type>(0));
ASSERT_EQ(view_test_h(10), static_cast<value_type>(0));
}
else if (name == "small-b") {
EXPECT_EQ(view_test_h(0), static_cast<value_type>(1));
EXPECT_EQ(view_test_h(1), static_cast<value_type>(2));
EXPECT_EQ(view_test_h(2), static_cast<value_type>(3));
EXPECT_EQ(view_test_h(3), static_cast<value_type>(4));
EXPECT_EQ(view_test_h(4), static_cast<value_type>(5));
EXPECT_EQ(view_test_h(5), static_cast<value_type>(6));
EXPECT_EQ(view_test_h(6), static_cast<value_type>(8));
EXPECT_EQ(view_test_h(7), static_cast<value_type>(9));
EXPECT_EQ(view_test_h(8), static_cast<value_type>(8));
EXPECT_EQ(view_test_h(9), static_cast<value_type>(0));
EXPECT_EQ(view_test_h(10), static_cast<value_type>(0));
EXPECT_EQ(view_test_h(11), static_cast<value_type>(0));
EXPECT_EQ(view_test_h(12), static_cast<value_type>(0));
ASSERT_EQ(view_test_h(0), static_cast<value_type>(1));
ASSERT_EQ(view_test_h(1), static_cast<value_type>(2));
ASSERT_EQ(view_test_h(2), static_cast<value_type>(3));
ASSERT_EQ(view_test_h(3), static_cast<value_type>(4));
ASSERT_EQ(view_test_h(4), static_cast<value_type>(5));
ASSERT_EQ(view_test_h(5), static_cast<value_type>(6));
ASSERT_EQ(view_test_h(6), static_cast<value_type>(8));
ASSERT_EQ(view_test_h(7), static_cast<value_type>(9));
ASSERT_EQ(view_test_h(8), static_cast<value_type>(8));
ASSERT_EQ(view_test_h(9), static_cast<value_type>(0));
ASSERT_EQ(view_test_h(10), static_cast<value_type>(0));
ASSERT_EQ(view_test_h(11), static_cast<value_type>(0));
ASSERT_EQ(view_test_h(12), static_cast<value_type>(0));
}
else if (name == "medium" || name == "large") {
@ -230,7 +230,7 @@ void verify_data(const std::string& name, ViewTypeFrom view_from,
(void)std_r;
for (std::size_t i = 0; i < view_from_h.extent(0); ++i) {
EXPECT_EQ(view_test_h(i), tmp[i]);
ASSERT_EQ(view_test_h(i), tmp[i]);
}
}
@ -273,7 +273,7 @@ void run_single_scenario(const InfoType& scenario_info, Args... args) {
KE::unique_copy(exespace(), KE::cbegin(view_from), KE::cend(view_from),
KE::begin(view_dest), args...);
verify_data(name, view_from, view_dest, args...);
EXPECT_EQ(rit, (KE::begin(view_dest) + n));
ASSERT_EQ(rit, (KE::begin(view_dest) + n));
}
{
@ -283,7 +283,7 @@ void run_single_scenario(const InfoType& scenario_info, Args... args) {
KE::unique_copy("label", exespace(), KE::cbegin(view_from),
KE::cend(view_from), KE::begin(view_dest), args...);
verify_data(name, view_from, view_dest, args...);
EXPECT_EQ(rit, (KE::begin(view_dest) + n));
ASSERT_EQ(rit, (KE::begin(view_dest) + n));
}
{
@ -291,7 +291,7 @@ void run_single_scenario(const InfoType& scenario_info, Args... args) {
create_view<ValueType>(Tag{}, view_ext, "unique_copy_dest");
auto rit = KE::unique_copy(exespace(), view_from, view_dest, args...);
verify_data(name, view_from, view_dest, args...);
EXPECT_EQ(rit, (KE::begin(view_dest) + n));
ASSERT_EQ(rit, (KE::begin(view_dest) + n));
}
{
@ -300,7 +300,7 @@ void run_single_scenario(const InfoType& scenario_info, Args... args) {
auto rit =
KE::unique_copy("label", exespace(), view_from, view_dest, args...);
verify_data(name, view_from, view_dest, args...);
EXPECT_EQ(rit, (KE::begin(view_dest) + n));
ASSERT_EQ(rit, (KE::begin(view_dest) + n));
}
Kokkos::fence();