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

Update Kokkos library in LAMMPS to v4.3.0

Browse Source
This commit is contained in:
Stan Gerald Moore
2024-04-05 08:20:57 -06:00
parent 2927471892
commit f382eac5c2
393 changed files with 13253 additions and 12499 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@ -30,5 +30,5 @@ KOKKOS_LIB_INCLUDE_DIRECTORIES(kokkosalgorithms
${CMAKE_CURRENT_SOURCE_DIR}
)
KOKKOS_LINK_TPL(kokkoscontainers PUBLIC ROCTHRUST)
KOKKOS_LINK_TPL(kokkoscore PUBLIC ONEDPL)

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

@ -849,18 +849,17 @@ class Random_XorShift64 {
return drand(end - start) + start;
}
// Marsaglia polar method for drawing a standard normal distributed random
// Box-muller method for drawing a standard normal distributed random
// number
KOKKOS_INLINE_FUNCTION
double normal() {
double S = 2.0;
double U;
while (S >= 1.0) {
U = 2.0 * drand() - 1.0;
const double V = 2.0 * drand() - 1.0;
S = U * U + V * V;
}
return U * std::sqrt(-2.0 * std::log(S) / S);
constexpr auto two_pi = 2 * Kokkos::numbers::pi_v<double>;
const double u = drand();
const double v = drand();
const double r = Kokkos::sqrt(-2.0 * Kokkos::log(u));
const double theta = v * two_pi;
return r * Kokkos::cos(theta);
}
KOKKOS_INLINE_FUNCTION
@ -1094,18 +1093,17 @@ class Random_XorShift1024 {
return drand(end - start) + start;
}
// Marsaglia polar method for drawing a standard normal distributed random
// Box-muller method for drawing a standard normal distributed random
// number
KOKKOS_INLINE_FUNCTION
double normal() {
double S = 2.0;
double U;
while (S >= 1.0) {
U = 2.0 * drand() - 1.0;
const double V = 2.0 * drand() - 1.0;
S = U * U + V * V;
}
return U * std::sqrt(-2.0 * std::log(S) / S);
constexpr auto two_pi = 2 * Kokkos::numbers::pi_v<double>;
const double u = drand();
const double v = drand();
const double r = Kokkos::sqrt(-2.0 * Kokkos::log(u));
const double theta = v * two_pi;
return r * Kokkos::cos(theta);
}
KOKKOS_INLINE_FUNCTION
@ -1545,13 +1543,23 @@ template <class ViewType, class RandomPool, class IndexType = int64_t>
void fill_random(ViewType a, RandomPool g,
typename ViewType::const_value_type begin,
typename ViewType::const_value_type end) {
fill_random(typename ViewType::execution_space{}, a, g, begin, end);
Kokkos::fence(
"fill_random: fence before since no execution space instance provided");
typename ViewType::execution_space exec;
fill_random(exec, a, g, begin, end);
exec.fence(
"fill_random: fence after since no execution space instance provided");
}
template <class ViewType, class RandomPool, class IndexType = int64_t>
void fill_random(ViewType a, RandomPool g,
typename ViewType::const_value_type range) {
fill_random(typename ViewType::execution_space{}, a, g, 0, range);
Kokkos::fence(
"fill_random: fence before since no execution space instance provided");
typename ViewType::execution_space exec;
fill_random(exec, a, g, 0, range);
exec.fence(
"fill_random: fence after since no execution space instance provided");
}
} // namespace Kokkos

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

@ -23,6 +23,7 @@
#include "sorting/Kokkos_BinSortPublicAPI.hpp"
#include "sorting/Kokkos_SortPublicAPI.hpp"
#include "sorting/Kokkos_SortByKeyPublicAPI.hpp"
#include "sorting/Kokkos_NestedSortPublicAPI.hpp"
#ifdef KOKKOS_IMPL_PUBLIC_INCLUDE_NOTDEFINED_SORT

1
lib/kokkos/algorithms/src/Kokkos_StdAlgorithms.hpp
View File

@ -35,7 +35,6 @@
// following the std classification.
// modifying ops
#include "std_algorithms/Kokkos_Swap.hpp"
#include "std_algorithms/Kokkos_IterSwap.hpp"
// non-modifying sequence

117
lib/kokkos/algorithms/src/sorting/Kokkos_SortByKeyPublicAPI.hpp Normal file
View File

@ -0,0 +1,117 @@
//@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_SORT_BY_KEY_PUBLIC_API_HPP_
#define KOKKOS_SORT_BY_KEY_PUBLIC_API_HPP_
#include "./impl/Kokkos_SortByKeyImpl.hpp"
#include <Kokkos_Core.hpp>
#include <algorithm>
namespace Kokkos::Experimental {
// ---------------------------------------------------------------
// basic overloads
// ---------------------------------------------------------------
template <class ExecutionSpace, class KeysDataType, class... KeysProperties,
class ValuesDataType, class... ValuesProperties>
void sort_by_key(
const ExecutionSpace& exec,
const Kokkos::View<KeysDataType, KeysProperties...>& keys,
const Kokkos::View<ValuesDataType, ValuesProperties...>& values) {
// constraints
using KeysType = Kokkos::View<KeysDataType, KeysProperties...>;
using ValuesType = Kokkos::View<ValuesDataType, ValuesProperties...>;
::Kokkos::Impl::static_assert_is_admissible_to_kokkos_sort_by_key(keys);
::Kokkos::Impl::static_assert_is_admissible_to_kokkos_sort_by_key(values);
static_assert(SpaceAccessibility<ExecutionSpace,
typename KeysType::memory_space>::accessible,
"Kokkos::sort: execution space instance is not able to access "
"the memory space of the keys View argument!");
static_assert(
SpaceAccessibility<ExecutionSpace,
typename ValuesType::memory_space>::accessible,
"Kokkos::sort: execution space instance is not able to access "
"the memory space of the values View argument!");
static_assert(KeysType::static_extent(0) == 0 ||
ValuesType::static_extent(0) == 0 ||
KeysType::static_extent(0) == ValuesType::static_extent(0));
if (values.size() != keys.size())
Kokkos::abort((std::string("values and keys extents must be the same. The "
"values extent is ") +
std::to_string(values.size()) + ", and the keys extent is " +
std::to_string(keys.size()) + ".")
.c_str());
if (keys.extent(0) <= 1) {
return;
}
::Kokkos::Impl::sort_by_key_device_view_without_comparator(exec, keys,
values);
}
// ---------------------------------------------------------------
// overloads supporting a custom comparator
// ---------------------------------------------------------------
template <class ExecutionSpace, class ComparatorType, class KeysDataType,
class... KeysProperties, class ValuesDataType,
class... ValuesProperties>
void sort_by_key(
const ExecutionSpace& exec,
const Kokkos::View<KeysDataType, KeysProperties...>& keys,
const Kokkos::View<ValuesDataType, ValuesProperties...>& values,
const ComparatorType& comparator) {
// constraints
using KeysType = Kokkos::View<KeysDataType, KeysProperties...>;
using ValuesType = Kokkos::View<ValuesDataType, ValuesProperties...>;
::Kokkos::Impl::static_assert_is_admissible_to_kokkos_sort_by_key(keys);
::Kokkos::Impl::static_assert_is_admissible_to_kokkos_sort_by_key(values);
static_assert(SpaceAccessibility<ExecutionSpace,
typename KeysType::memory_space>::accessible,
"Kokkos::sort: execution space instance is not able to access "
"the memory space of the keys View argument!");
static_assert(
SpaceAccessibility<ExecutionSpace,
typename ValuesType::memory_space>::accessible,
"Kokkos::sort: execution space instance is not able to access "
"the memory space of the values View argument!");
static_assert(KeysType::static_extent(0) == 0 ||
ValuesType::static_extent(0) == 0 ||
KeysType::static_extent(0) == ValuesType::static_extent(0));
if (values.size() != keys.size())
Kokkos::abort((std::string("values and keys extents must be the same. The "
"values extent is ") +
std::to_string(values.size()) + ", and the keys extent is " +
std::to_string(keys.size()) + ".")
.c_str());
if (keys.extent(0) <= 1) {
return;
}
::Kokkos::Impl::sort_by_key_device_view_with_comparator(exec, keys, values,
comparator);
}
} // namespace Kokkos::Experimental
#endif

6
lib/kokkos/algorithms/src/sorting/Kokkos_SortPublicAPI.hpp
View File

@ -29,7 +29,7 @@ namespace Kokkos {
// ---------------------------------------------------------------
template <class ExecutionSpace, class DataType, class... Properties>
void sort([[maybe_unused]] const ExecutionSpace& exec,
void sort(const ExecutionSpace& exec,
const Kokkos::View<DataType, Properties...>& view) {
// constraints
using ViewType = Kokkos::View<DataType, Properties...>;
@ -52,6 +52,7 @@ void sort([[maybe_unused]] const ExecutionSpace& exec,
}
if constexpr (Impl::better_off_calling_std_sort_v<ExecutionSpace>) {
exec.fence("Kokkos::sort without comparator use std::sort");
auto first = ::Kokkos::Experimental::begin(view);
auto last = ::Kokkos::Experimental::end(view);
std::sort(first, last);
@ -82,7 +83,7 @@ void sort(const Kokkos::View<DataType, Properties...>& view) {
// ---------------------------------------------------------------
template <class ExecutionSpace, class ComparatorType, class DataType,
class... Properties>
void sort([[maybe_unused]] const ExecutionSpace& exec,
void sort(const ExecutionSpace& exec,
const Kokkos::View<DataType, Properties...>& view,
const ComparatorType& comparator) {
// constraints
@ -105,6 +106,7 @@ void sort([[maybe_unused]] const ExecutionSpace& exec,
}
if constexpr (Impl::better_off_calling_std_sort_v<ExecutionSpace>) {
exec.fence("Kokkos::sort with comparator use std::sort");
auto first = ::Kokkos::Experimental::begin(view);
auto last = ::Kokkos::Experimental::end(view);
std::sort(first, last, comparator);

3
lib/kokkos/algorithms/src/sorting/impl/Kokkos_NestedSortImpl.hpp
View File

@ -18,7 +18,6 @@
#define KOKKOS_NESTED_SORT_IMPL_HPP_
#include <Kokkos_Core.hpp>
#include <std_algorithms/Kokkos_Swap.hpp>
namespace Kokkos {
namespace Experimental {
@ -99,7 +98,7 @@ KOKKOS_INLINE_FUNCTION void sort_nested_impl(
keyView(elem1) = key2;
keyView(elem2) = key1;
if constexpr (!std::is_same_v<ValueViewType, std::nullptr_t>) {
Kokkos::Experimental::swap(valueView(elem1), valueView(elem2));
Kokkos::kokkos_swap(valueView(elem1), valueView(elem2));
}
}
}

401
lib/kokkos/algorithms/src/sorting/impl/Kokkos_SortByKeyImpl.hpp Normal file
View File

@ -0,0 +1,401 @@
//@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_SORT_BY_KEY_FREE_FUNCS_IMPL_HPP_
#define KOKKOS_SORT_BY_KEY_FREE_FUNCS_IMPL_HPP_
#include <Kokkos_Core.hpp>
#if defined(KOKKOS_ENABLE_CUDA)
// Workaround for `Instruction 'shfl' without '.sync' is not supported on
// .target sm_70 and higher from PTX ISA version 6.4`.
// Also see https://github.com/NVIDIA/cub/pull/170.
#if !defined(CUB_USE_COOPERATIVE_GROUPS)
#define CUB_USE_COOPERATIVE_GROUPS
#endif
#pragma GCC diagnostic push
#pragma GCC diagnostic ignored "-Wshadow"
#if defined(KOKKOS_COMPILER_CLANG)
// Some versions of Clang fail to compile Thrust, failing with errors like
// this:
// <snip>/thrust/system/cuda/detail/core/agent_launcher.h:557:11:
// error: use of undeclared identifier 'va_printf'
// The exact combination of versions for Clang and Thrust (or CUDA) for this
// failure was not investigated, however even very recent version combination
// (Clang 10.0.0 and Cuda 10.0) demonstrated failure.
//
// Defining _CubLog here locally allows us to avoid that code path, however
// disabling some debugging diagnostics
#pragma push_macro("_CubLog")
#ifdef _CubLog
#undef _CubLog
#endif
#define _CubLog
#include <thrust/device_ptr.h>
#include <thrust/sort.h>
#pragma pop_macro("_CubLog")
#else
#include <thrust/device_ptr.h>
#include <thrust/sort.h>
#endif
#pragma GCC diagnostic pop
#endif
#if defined(KOKKOS_ENABLE_ROCTHRUST)
#include <thrust/device_ptr.h>
#include <thrust/sort.h>
#endif
#if defined(KOKKOS_ENABLE_ONEDPL) && \
(ONEDPL_VERSION_MAJOR > 2022 || \
(ONEDPL_VERSION_MAJOR == 2022 && ONEDPL_VERSION_MINOR >= 2))
#define KOKKOS_ONEDPL_HAS_SORT_BY_KEY
#include <oneapi/dpl/execution>
#include <oneapi/dpl/algorithm>
#endif
namespace Kokkos::Impl {
template <typename T>
constexpr inline bool is_admissible_to_kokkos_sort_by_key =
::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,
Kokkos::LayoutRight>::value ||
std::is_same<typename T::traits::array_layout,
Kokkos::LayoutStride>::value);
template <class ViewType>
KOKKOS_INLINE_FUNCTION constexpr void
static_assert_is_admissible_to_kokkos_sort_by_key(const ViewType& /* view */) {
static_assert(is_admissible_to_kokkos_sort_by_key<ViewType>,
"Kokkos::sort_by_key only accepts 1D values View with "
"LayoutRight, LayoutLeft or LayoutStride.");
}
// For the fallback implementation for sort_by_key using Kokkos::sort, we need
// to consider if Kokkos::sort defers to the fallback implementation that copies
// the array to the host and uses std::sort, see
// copy_to_host_run_stdsort_copy_back() in impl/Kokkos_SortImpl.hpp. If
// sort_on_device_v is true, we assume that std::sort doesn't copy data.
// Otherwise, we manually copy all data to the host and provide Kokkos::sort
// with a host execution space.
template <class ExecutionSpace, class Layout>
inline constexpr bool sort_on_device_v = false;
#if defined(KOKKOS_ENABLE_CUDA)
template <class Layout>
inline constexpr bool sort_on_device_v<Kokkos::Cuda, Layout> = true;
template <class KeysDataType, class... KeysProperties, class ValuesDataType,
class... ValuesProperties, class... MaybeComparator>
void sort_by_key_cudathrust(
const Kokkos::Cuda& exec,
const Kokkos::View<KeysDataType, KeysProperties...>& keys,
const Kokkos::View<ValuesDataType, ValuesProperties...>& values,
MaybeComparator&&... maybeComparator) {
const auto policy = thrust::cuda::par.on(exec.cuda_stream());
auto keys_first = ::Kokkos::Experimental::begin(keys);
auto keys_last = ::Kokkos::Experimental::end(keys);
auto values_first = ::Kokkos::Experimental::begin(values);
thrust::sort_by_key(policy, keys_first, keys_last, values_first,
std::forward<MaybeComparator>(maybeComparator)...);
}
#endif
#if defined(KOKKOS_ENABLE_ROCTHRUST)
template <class Layout>
inline constexpr bool sort_on_device_v<Kokkos::HIP, Layout> = true;
template <class KeysDataType, class... KeysProperties, class ValuesDataType,
class... ValuesProperties, class... MaybeComparator>
void sort_by_key_rocthrust(
const Kokkos::HIP& exec,
const Kokkos::View<KeysDataType, KeysProperties...>& keys,
const Kokkos::View<ValuesDataType, ValuesProperties...>& values,
MaybeComparator&&... maybeComparator) {
const auto policy = thrust::hip::par.on(exec.hip_stream());
auto keys_first = ::Kokkos::Experimental::begin(keys);
auto keys_last = ::Kokkos::Experimental::end(keys);
auto values_first = ::Kokkos::Experimental::begin(values);
thrust::sort_by_key(policy, keys_first, keys_last, values_first,
std::forward<MaybeComparator>(maybeComparator)...);
}
#endif
#if defined(KOKKOS_ENABLE_ONEDPL)
template <class Layout>
inline constexpr bool sort_on_device_v<Kokkos::Experimental::SYCL, Layout> =
std::is_same_v<Layout, Kokkos::LayoutLeft> ||
std::is_same_v<Layout, Kokkos::LayoutRight>;
#ifdef KOKKOS_ONEDPL_HAS_SORT_BY_KEY
template <class KeysDataType, class... KeysProperties, class ValuesDataType,
class... ValuesProperties, class... MaybeComparator>
void sort_by_key_onedpl(
const Kokkos::Experimental::SYCL& exec,
const Kokkos::View<KeysDataType, KeysProperties...>& keys,
const Kokkos::View<ValuesDataType, ValuesProperties...>& values,
MaybeComparator&&... maybeComparator) {
if (keys.stride(0) != 1 && values.stride(0) != 1) {
Kokkos::abort(
"SYCL sort_by_key only supports rank-1 Views with stride(0) = 1.");
}
// Can't use Experimental::begin/end here since the oneDPL then assumes that
// the data is on the host.
auto queue = exec.sycl_queue();
auto policy = oneapi::dpl::execution::make_device_policy(queue);
const int n = keys.extent(0);
oneapi::dpl::sort_by_key(policy, keys.data(), keys.data() + n, values.data(),
std::forward<MaybeComparator>(maybeComparator)...);
}
#endif
#endif
template <typename ExecutionSpace, typename PermutationView, typename ViewType>
void applyPermutation(const ExecutionSpace& space,
const PermutationView& permutation,
const ViewType& view) {
static_assert(std::is_integral<typename PermutationView::value_type>::value);
auto view_copy = Kokkos::create_mirror(
Kokkos::view_alloc(space, typename ExecutionSpace::memory_space{},
Kokkos::WithoutInitializing),
view);
Kokkos::deep_copy(space, view_copy, view);
Kokkos::parallel_for(
"Kokkos::sort_by_key_via_sort::permute_" + view.label(),
Kokkos::RangePolicy<ExecutionSpace>(space, 0, view.extent(0)),
KOKKOS_LAMBDA(int i) { view(i) = view_copy(permutation(i)); });
}
template <class ExecutionSpace, class KeysDataType, class... KeysProperties,
class ValuesDataType, class... ValuesProperties,
class... MaybeComparator>
void sort_by_key_via_sort(
const ExecutionSpace& exec,
const Kokkos::View<KeysDataType, KeysProperties...>& keys,
const Kokkos::View<ValuesDataType, ValuesProperties...>& values,
MaybeComparator&&... maybeComparator) {
static_assert(sizeof...(MaybeComparator) <= 1);
auto const n = keys.size();
Kokkos::View<unsigned int*, ExecutionSpace> permute(
Kokkos::view_alloc(exec, Kokkos::WithoutInitializing,
"Kokkos::sort_by_key_via_sort::permute"),
n);
// iota
Kokkos::parallel_for(
"Kokkos::sort_by_key_via_sort::iota",
Kokkos::RangePolicy<ExecutionSpace>(exec, 0, n),
KOKKOS_LAMBDA(int i) { permute(i) = i; });
using Layout =
typename Kokkos::View<unsigned int*, ExecutionSpace>::array_layout;
if constexpr (!sort_on_device_v<ExecutionSpace, Layout>) {
auto host_keys = Kokkos::create_mirror_view(
Kokkos::view_alloc(Kokkos::HostSpace{}, Kokkos::WithoutInitializing),
keys);
auto host_permute = Kokkos::create_mirror_view(
Kokkos::view_alloc(Kokkos::HostSpace{}, Kokkos::WithoutInitializing),
permute);
Kokkos::deep_copy(exec, host_keys, keys);
Kokkos::deep_copy(exec, host_permute, permute);
exec.fence("Kokkos::Impl::sort_by_key_via_sort: before host sort");
Kokkos::DefaultHostExecutionSpace host_exec;
if constexpr (sizeof...(MaybeComparator) == 0) {
Kokkos::sort(
host_exec, host_permute,
KOKKOS_LAMBDA(int i, int j) { return host_keys(i) < host_keys(j); });
} else {
auto keys_comparator =
std::get<0>(std::tuple<MaybeComparator...>(maybeComparator...));
Kokkos::sort(
host_exec, host_permute, KOKKOS_LAMBDA(int i, int j) {
return keys_comparator(host_keys(i), host_keys(j));
});
}
host_exec.fence("Kokkos::Impl::sort_by_key_via_sort: after host sort");
Kokkos::deep_copy(exec, permute, host_permute);
} else {
#ifdef KOKKOS_ENABLE_SYCL
auto* raw_keys_in_comparator = keys.data();
auto stride = keys.stride(0);
if constexpr (sizeof...(MaybeComparator) == 0) {
Kokkos::sort(
exec, permute, KOKKOS_LAMBDA(int i, int j) {
return raw_keys_in_comparator[i * stride] <
raw_keys_in_comparator[j * stride];
});
} else {
auto keys_comparator =
std::get<0>(std::tuple<MaybeComparator...>(maybeComparator...));
Kokkos::sort(
exec, permute, KOKKOS_LAMBDA(int i, int j) {
return keys_comparator(raw_keys_in_comparator[i * stride],
raw_keys_in_comparator[j * stride]);
});
}
#else
if constexpr (sizeof...(MaybeComparator) == 0) {
Kokkos::sort(
exec, permute,
KOKKOS_LAMBDA(int i, int j) { return keys(i) < keys(j); });
} else {
auto keys_comparator =
std::get<0>(std::tuple<MaybeComparator...>(maybeComparator...));
Kokkos::sort(
exec, permute, KOKKOS_LAMBDA(int i, int j) {
return keys_comparator(keys(i), keys(j));
});
}
#endif
}
applyPermutation(exec, permute, keys);
applyPermutation(exec, permute, values);
}
// ------------------------------------------------------
//
// specialize cases for sorting by key without comparator
//
// ------------------------------------------------------
#if defined(KOKKOS_ENABLE_CUDA)
template <class KeysDataType, class... KeysProperties, class ValuesDataType,
class... ValuesProperties>
void sort_by_key_device_view_without_comparator(
const Kokkos::Cuda& exec,
const Kokkos::View<KeysDataType, KeysProperties...>& keys,
const Kokkos::View<ValuesDataType, ValuesProperties...>& values) {
sort_by_key_cudathrust(exec, keys, values);
}
#endif
#if defined(KOKKOS_ENABLE_ROCTHRUST)
template <class KeysDataType, class... KeysProperties, class ValuesDataType,
class... ValuesProperties>
void sort_by_key_device_view_without_comparator(
const Kokkos::HIP& exec,
const Kokkos::View<KeysDataType, KeysProperties...>& keys,
const Kokkos::View<ValuesDataType, ValuesProperties...>& values) {
sort_by_key_rocthrust(exec, keys, values);
}
#endif
#if defined(KOKKOS_ENABLE_ONEDPL)
template <class KeysDataType, class... KeysProperties, class ValuesDataType,
class... ValuesProperties>
void sort_by_key_device_view_without_comparator(
const Kokkos::Experimental::SYCL& exec,
const Kokkos::View<KeysDataType, KeysProperties...>& keys,
const Kokkos::View<ValuesDataType, ValuesProperties...>& values) {
#ifdef KOKKOS_ONEDPL_HAS_SORT_BY_KEY
if (keys.stride(0) == 1 && values.stride(0) == 1)
sort_by_key_onedpl(exec, keys, values);
else
#endif
sort_by_key_via_sort(exec, keys, values);
}
#endif
// fallback case
template <class ExecutionSpace, class KeysDataType, class... KeysProperties,
class ValuesDataType, class... ValuesProperties>
std::enable_if_t<Kokkos::is_execution_space<ExecutionSpace>::value>
sort_by_key_device_view_without_comparator(
const ExecutionSpace& exec,
const Kokkos::View<KeysDataType, KeysProperties...>& keys,
const Kokkos::View<ValuesDataType, ValuesProperties...>& values) {
sort_by_key_via_sort(exec, keys, values);
}
// ---------------------------------------------------
//
// specialize cases for sorting by key with comparator
//
// ---------------------------------------------------
#if defined(KOKKOS_ENABLE_CUDA)
template <class ComparatorType, class KeysDataType, class... KeysProperties,
class ValuesDataType, class... ValuesProperties>
void sort_by_key_device_view_with_comparator(
const Kokkos::Cuda& exec,
const Kokkos::View<KeysDataType, KeysProperties...>& keys,
const Kokkos::View<ValuesDataType, ValuesProperties...>& values,
const ComparatorType& comparator) {
sort_by_key_cudathrust(exec, keys, values, comparator);
}
#endif
#if defined(KOKKOS_ENABLE_ROCTHRUST)
template <class ComparatorType, class KeysDataType, class... KeysProperties,
class ValuesDataType, class... ValuesProperties>
void sort_by_key_device_view_with_comparator(
const Kokkos::HIP& exec,
const Kokkos::View<KeysDataType, KeysProperties...>& keys,
const Kokkos::View<ValuesDataType, ValuesProperties...>& values,
const ComparatorType& comparator) {
sort_by_key_rocthrust(exec, keys, values, comparator);
}
#endif
#if defined(KOKKOS_ENABLE_ONEDPL)
template <class ComparatorType, class KeysDataType, class... KeysProperties,
class ValuesDataType, class... ValuesProperties>
void sort_by_key_device_view_with_comparator(
const Kokkos::Experimental::SYCL& exec,
const Kokkos::View<KeysDataType, KeysProperties...>& keys,
const Kokkos::View<ValuesDataType, ValuesProperties...>& values,
const ComparatorType& comparator) {
#ifdef KOKKOS_ONEDPL_HAS_SORT_BY_KEY
if (keys.stride(0) == 1 && values.stride(0) == 1)
sort_by_key_onedpl(exec, keys, values, comparator);
else
#endif
sort_by_key_via_sort(exec, keys, values, comparator);
}
#endif
// fallback case
template <class ComparatorType, class ExecutionSpace, class KeysDataType,
class... KeysProperties, class ValuesDataType,
class... ValuesProperties>
std::enable_if_t<Kokkos::is_execution_space<ExecutionSpace>::value>
sort_by_key_device_view_with_comparator(
const ExecutionSpace& exec,
const Kokkos::View<KeysDataType, KeysProperties...>& keys,
const Kokkos::View<ValuesDataType, ValuesProperties...>& values,
const ComparatorType& comparator) {
sort_by_key_via_sort(exec, keys, values, comparator);
}
#undef KOKKOS_ONEDPL_HAS_SORT_BY_KEY
} // namespace Kokkos::Impl
#endif

42
lib/kokkos/algorithms/src/sorting/impl/Kokkos_SortImpl.hpp
View File

@ -63,6 +63,11 @@
#endif
#if defined(KOKKOS_ENABLE_ROCTHRUST)
#include <thrust/device_ptr.h>
#include <thrust/sort.h>
#endif
#if defined(KOKKOS_ENABLE_ONEDPL)
#include <oneapi/dpl/execution>
#include <oneapi/dpl/algorithm>
@ -184,6 +189,26 @@ void sort_cudathrust(const Cuda& space,
}
#endif
#if defined(KOKKOS_ENABLE_ROCTHRUST)
template <class DataType, class... Properties, class... MaybeComparator>
void sort_rocthrust(const HIP& space,
const Kokkos::View<DataType, Properties...>& view,
MaybeComparator&&... maybeComparator) {
using ViewType = Kokkos::View<DataType, Properties...>;
static_assert(ViewType::rank == 1,
"Kokkos::sort: currently only supports rank-1 Views.");
if (view.extent(0) <= 1) {
return;
}
const auto exec = thrust::hip::par.on(space.hip_stream());
auto first = ::Kokkos::Experimental::begin(view);
auto last = ::Kokkos::Experimental::end(view);
thrust::sort(exec, first, last,
std::forward<MaybeComparator>(maybeComparator)...);
}
#endif
#if defined(KOKKOS_ENABLE_ONEDPL)
template <class DataType, class... Properties, class... MaybeComparator>
void sort_onedpl(const Kokkos::Experimental::SYCL& space,
@ -274,6 +299,14 @@ void sort_device_view_without_comparator(
}
#endif
#if defined(KOKKOS_ENABLE_ROCTHRUST)
template <class DataType, class... Properties>
void sort_device_view_without_comparator(
const HIP& exec, const Kokkos::View<DataType, Properties...>& view) {
sort_rocthrust(exec, view);
}
#endif
#if defined(KOKKOS_ENABLE_ONEDPL)
template <class DataType, class... Properties>
void sort_device_view_without_comparator(
@ -320,6 +353,15 @@ void sort_device_view_with_comparator(
}
#endif
#if defined(KOKKOS_ENABLE_ROCTHRUST)
template <class ComparatorType, class DataType, class... Properties>
void sort_device_view_with_comparator(
const HIP& exec, const Kokkos::View<DataType, Properties...>& view,
const ComparatorType& comparator) {
sort_rocthrust(exec, view, comparator);
}
#endif
#if defined(KOKKOS_ENABLE_ONEDPL)
template <class ComparatorType, class DataType, class... Properties>
void sort_device_view_with_comparator(

6
lib/kokkos/algorithms/src/std_algorithms/Kokkos_Copy.hpp
View File

@ -50,7 +50,7 @@ template <
std::enable_if_t<::Kokkos::is_execution_space_v<ExecutionSpace>, int> = 0>
auto copy(const ExecutionSpace& ex,
const ::Kokkos::View<DataType1, Properties1...>& source,
::Kokkos::View<DataType2, Properties2...>& dest) {
const ::Kokkos::View<DataType2, Properties2...>& dest) {
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(source);
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(dest);
@ -66,7 +66,7 @@ template <
std::enable_if_t<::Kokkos::is_execution_space_v<ExecutionSpace>, int> = 0>
auto copy(const std::string& label, const ExecutionSpace& ex,
const ::Kokkos::View<DataType1, Properties1...>& source,
::Kokkos::View<DataType2, Properties2...>& dest) {
const ::Kokkos::View<DataType2, Properties2...>& dest) {
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(source);
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(dest);
@ -93,7 +93,7 @@ template <typename TeamHandleType, typename DataType1, typename... Properties1,
KOKKOS_FUNCTION auto copy(
const TeamHandleType& teamHandle,
const ::Kokkos::View<DataType1, Properties1...>& source,
::Kokkos::View<DataType2, Properties2...>& dest) {
const ::Kokkos::View<DataType2, Properties2...>& dest) {
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(source);
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(dest);

6
lib/kokkos/algorithms/src/std_algorithms/Kokkos_CopyBackward.hpp
View File

@ -50,7 +50,7 @@ template <
std::enable_if_t<::Kokkos::is_execution_space_v<ExecutionSpace>, int> = 0>
auto copy_backward(const ExecutionSpace& ex,
const ::Kokkos::View<DataType1, Properties1...>& source,
::Kokkos::View<DataType2, Properties2...>& dest) {
const ::Kokkos::View<DataType2, Properties2...>& dest) {
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(source);
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(dest);
@ -65,7 +65,7 @@ template <
std::enable_if_t<::Kokkos::is_execution_space_v<ExecutionSpace>, int> = 0>
auto copy_backward(const std::string& label, const ExecutionSpace& ex,
const ::Kokkos::View<DataType1, Properties1...>& source,
::Kokkos::View<DataType2, Properties2...>& dest) {
const ::Kokkos::View<DataType2, Properties2...>& dest) {
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(source);
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(dest);
@ -92,7 +92,7 @@ template <typename TeamHandleType, typename DataType1, typename... Properties1,
KOKKOS_FUNCTION auto copy_backward(
const TeamHandleType& teamHandle,
const ::Kokkos::View<DataType1, Properties1...>& source,
::Kokkos::View<DataType2, Properties2...>& dest) {
const ::Kokkos::View<DataType2, Properties2...>& dest) {
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(source);
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(dest);

8
lib/kokkos/algorithms/src/std_algorithms/Kokkos_CopyIf.hpp
View File

@ -54,7 +54,8 @@ template <
std::enable_if_t<::Kokkos::is_execution_space_v<ExecutionSpace>, int> = 0>
auto copy_if(const ExecutionSpace& ex,
const ::Kokkos::View<DataType1, Properties1...>& source,
::Kokkos::View<DataType2, Properties2...>& dest, Predicate pred) {
const ::Kokkos::View<DataType2, Properties2...>& dest,
Predicate pred) {
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(source);
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(dest);
@ -69,7 +70,8 @@ template <
std::enable_if_t<::Kokkos::is_execution_space_v<ExecutionSpace>, int> = 0>
auto copy_if(const std::string& label, const ExecutionSpace& ex,
const ::Kokkos::View<DataType1, Properties1...>& source,
::Kokkos::View<DataType2, Properties2...>& dest, Predicate pred) {
const ::Kokkos::View<DataType2, Properties2...>& dest,
Predicate pred) {
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(source);
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(dest);
@ -96,7 +98,7 @@ template <typename TeamHandleType, typename DataType1, typename... Properties1,
KOKKOS_FUNCTION auto copy_if(
const TeamHandleType& teamHandle,
const ::Kokkos::View<DataType1, Properties1...>& source,
::Kokkos::View<DataType2, Properties2...>& dest, Predicate pred) {
const ::Kokkos::View<DataType2, Properties2...>& dest, Predicate pred) {
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(source);
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(dest);

6
lib/kokkos/algorithms/src/std_algorithms/Kokkos_CopyN.hpp
View File

@ -51,7 +51,7 @@ template <
std::enable_if_t<::Kokkos::is_execution_space_v<ExecutionSpace>, int> = 0>
auto copy_n(const ExecutionSpace& ex,
const ::Kokkos::View<DataType1, Properties1...>& source, Size count,
::Kokkos::View<DataType2, Properties2...>& dest) {
const ::Kokkos::View<DataType2, Properties2...>& dest) {
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(source);
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(dest);
@ -66,7 +66,7 @@ template <
std::enable_if_t<::Kokkos::is_execution_space_v<ExecutionSpace>, int> = 0>
auto copy_n(const std::string& label, const ExecutionSpace& ex,
const ::Kokkos::View<DataType1, Properties1...>& source, Size count,
::Kokkos::View<DataType2, Properties2...>& dest) {
const ::Kokkos::View<DataType2, Properties2...>& dest) {
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(source);
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(dest);
@ -93,7 +93,7 @@ template <typename TeamHandleType, typename DataType1, typename... Properties1,
KOKKOS_FUNCTION auto copy_n(
const TeamHandleType& teamHandle,
const ::Kokkos::View<DataType1, Properties1...>& source, Size count,
::Kokkos::View<DataType2, Properties2...>& dest) {
const ::Kokkos::View<DataType2, Properties2...>& dest) {
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(source);
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(dest);

12
lib/kokkos/algorithms/src/std_algorithms/Kokkos_Equal.hpp
View File

@ -80,7 +80,7 @@ template <
std::enable_if_t<::Kokkos::is_execution_space_v<ExecutionSpace>, int> = 0>
bool equal(const ExecutionSpace& ex,
const ::Kokkos::View<DataType1, Properties1...>& view1,
::Kokkos::View<DataType2, Properties2...>& view2) {
const ::Kokkos::View<DataType2, Properties2...>& view2) {
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(view1);
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(view2);
@ -96,7 +96,7 @@ template <
std::enable_if_t<::Kokkos::is_execution_space_v<ExecutionSpace>, int> = 0>
bool equal(const std::string& label, const ExecutionSpace& ex,
const ::Kokkos::View<DataType1, Properties1...>& view1,
::Kokkos::View<DataType2, Properties2...>& view2) {
const ::Kokkos::View<DataType2, Properties2...>& view2) {
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(view1);
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(view2);
@ -111,7 +111,7 @@ template <
std::enable_if_t<::Kokkos::is_execution_space_v<ExecutionSpace>, int> = 0>
bool equal(const ExecutionSpace& ex,
const ::Kokkos::View<DataType1, Properties1...>& view1,
::Kokkos::View<DataType2, Properties2...>& view2,
const ::Kokkos::View<DataType2, Properties2...>& view2,
BinaryPredicateType predicate) {
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(view1);
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(view2);
@ -128,7 +128,7 @@ template <
std::enable_if_t<::Kokkos::is_execution_space_v<ExecutionSpace>, int> = 0>
bool equal(const std::string& label, const ExecutionSpace& ex,
const ::Kokkos::View<DataType1, Properties1...>& view1,
::Kokkos::View<DataType2, Properties2...>& view2,
const ::Kokkos::View<DataType2, Properties2...>& view2,
BinaryPredicateType predicate) {
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(view1);
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(view2);
@ -227,7 +227,7 @@ template <typename TeamHandleType, typename DataType1, typename... Properties1,
KOKKOS_FUNCTION bool equal(
const TeamHandleType& teamHandle,
const ::Kokkos::View<DataType1, Properties1...>& view1,
::Kokkos::View<DataType2, Properties2...>& view2) {
const ::Kokkos::View<DataType2, Properties2...>& view2) {
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(view1);
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(view2);
@ -243,7 +243,7 @@ template <typename TeamHandleType, typename DataType1, typename... Properties1,
KOKKOS_FUNCTION bool equal(
const TeamHandleType& teamHandle,
const ::Kokkos::View<DataType1, Properties1...>& view1,
::Kokkos::View<DataType2, Properties2...>& view2,
const ::Kokkos::View<DataType2, Properties2...>& view2,
BinaryPredicateType predicate) {
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(view1);
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(view2);

13
lib/kokkos/algorithms/src/std_algorithms/Kokkos_IterSwap.hpp
View File

@ -19,7 +19,6 @@
#include <Kokkos_Core.hpp>
#include "impl/Kokkos_Constraints.hpp"
#include "Kokkos_Swap.hpp"
namespace Kokkos {
namespace Experimental {
@ -33,7 +32,7 @@ struct StdIterSwapFunctor {
KOKKOS_FUNCTION
void operator()(int i) const {
(void)i;
::Kokkos::Experimental::swap(*m_a, *m_b);
::Kokkos::kokkos_swap(*m_a, *m_b);
}
KOKKOS_FUNCTION
@ -58,6 +57,16 @@ void iter_swap(IteratorType1 a, IteratorType2 b) {
Impl::iter_swap_impl(a, b);
}
#ifdef KOKKOS_ENABLE_DEPRECATED_CODE_4
template <class T>
KOKKOS_DEPRECATED_WITH_COMMENT("Use Kokkos::kokkos_swap instead!")
KOKKOS_FUNCTION
void swap(T& a, T& b) noexcept(::Kokkos::kokkos_swap(std::declval<T&>(),
std::declval<T&>())) {
::Kokkos::kokkos_swap(a, b);
}
#endif
} // namespace Experimental
} // namespace Kokkos

15
lib/kokkos/algorithms/src/std_algorithms/Kokkos_LexicographicalCompare.hpp
View File

@ -54,7 +54,7 @@ template <
bool lexicographical_compare(
const ExecutionSpace& ex,
const ::Kokkos::View<DataType1, Properties1...>& view1,
::Kokkos::View<DataType2, Properties2...>& view2) {
const ::Kokkos::View<DataType2, Properties2...>& view2) {
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(view1);
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(view2);
@ -71,7 +71,7 @@ template <
bool lexicographical_compare(
const std::string& label, const ExecutionSpace& ex,
const ::Kokkos::View<DataType1, Properties1...>& view1,
::Kokkos::View<DataType2, Properties2...>& view2) {
const ::Kokkos::View<DataType2, Properties2...>& view2) {
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(view1);
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(view2);
@ -112,7 +112,8 @@ template <
bool lexicographical_compare(
const ExecutionSpace& ex,
const ::Kokkos::View<DataType1, Properties1...>& view1,
::Kokkos::View<DataType2, Properties2...>& view2, ComparatorType comp) {
const ::Kokkos::View<DataType2, Properties2...>& view2,
ComparatorType comp) {
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(view1);
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(view2);
@ -129,7 +130,8 @@ template <
bool lexicographical_compare(
const std::string& label, const ExecutionSpace& ex,
const ::Kokkos::View<DataType1, Properties1...>& view1,
::Kokkos::View<DataType2, Properties2...>& view2, ComparatorType comp) {
const ::Kokkos::View<DataType2, Properties2...>& view2,
ComparatorType comp) {
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(view1);
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(view2);
@ -161,7 +163,7 @@ template <class TeamHandleType, class DataType1, class... Properties1,
KOKKOS_FUNCTION bool lexicographical_compare(
const TeamHandleType& teamHandle,
const ::Kokkos::View<DataType1, Properties1...>& view1,
::Kokkos::View<DataType2, Properties2...>& view2) {
const ::Kokkos::View<DataType2, Properties2...>& view2) {
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(view1);
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(view2);
@ -187,7 +189,8 @@ template <class TeamHandleType, class DataType1, class... Properties1,
KOKKOS_FUNCTION bool lexicographical_compare(
const TeamHandleType& teamHandle,
const ::Kokkos::View<DataType1, Properties1...>& view1,
::Kokkos::View<DataType2, Properties2...>& view2, ComparatorType comp) {
const ::Kokkos::View<DataType2, Properties2...>& view2,
ComparatorType comp) {
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(view1);
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(view2);

6
lib/kokkos/algorithms/src/std_algorithms/Kokkos_Move.hpp
View File

@ -50,7 +50,7 @@ template <
std::enable_if_t<::Kokkos::is_execution_space_v<ExecutionSpace>, int> = 0>
auto move(const ExecutionSpace& ex,
const ::Kokkos::View<DataType1, Properties1...>& source,
::Kokkos::View<DataType2, Properties2...>& dest) {
const ::Kokkos::View<DataType2, Properties2...>& dest) {
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(source);
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(dest);
@ -64,7 +64,7 @@ template <
std::enable_if_t<::Kokkos::is_execution_space_v<ExecutionSpace>, int> = 0>
auto move(const std::string& label, const ExecutionSpace& ex,
const ::Kokkos::View<DataType1, Properties1...>& source,
::Kokkos::View<DataType2, Properties2...>& dest) {
const ::Kokkos::View<DataType2, Properties2...>& dest) {
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(source);
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(dest);
@ -92,7 +92,7 @@ template <typename TeamHandleType, typename DataType1, typename... Properties1,
KOKKOS_FUNCTION auto move(
const TeamHandleType& teamHandle,
const ::Kokkos::View<DataType1, Properties1...>& source,
::Kokkos::View<DataType2, Properties2...>& dest) {
const ::Kokkos::View<DataType2, Properties2...>& dest) {
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(source);
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(dest);

6
lib/kokkos/algorithms/src/std_algorithms/Kokkos_MoveBackward.hpp
View File

@ -41,7 +41,7 @@ template <
std::enable_if_t<::Kokkos::is_execution_space_v<ExecutionSpace>, int> = 0>
auto move_backward(const ExecutionSpace& ex,
const ::Kokkos::View<DataType1, Properties1...>& source,
::Kokkos::View<DataType2, Properties2...>& dest) {
const ::Kokkos::View<DataType2, Properties2...>& dest) {
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(source);
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(dest);
@ -65,7 +65,7 @@ template <
std::enable_if_t<::Kokkos::is_execution_space_v<ExecutionSpace>, int> = 0>
auto move_backward(const std::string& label, const ExecutionSpace& ex,
const ::Kokkos::View<DataType1, Properties1...>& source,
::Kokkos::View<DataType2, Properties2...>& dest) {
const ::Kokkos::View<DataType2, Properties2...>& dest) {
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(source);
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(dest);
@ -94,7 +94,7 @@ template <typename TeamHandleType, typename DataType1, typename... Properties1,
KOKKOS_FUNCTION auto move_backward(
const TeamHandleType& teamHandle,
const ::Kokkos::View<DataType1, Properties1...>& source,
::Kokkos::View<DataType2, Properties2...>& dest) {
const ::Kokkos::View<DataType2, Properties2...>& dest) {
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(source);
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(dest);

6
lib/kokkos/algorithms/src/std_algorithms/Kokkos_ReverseCopy.hpp
View File

@ -50,7 +50,7 @@ template <
std::enable_if_t<::Kokkos::is_execution_space_v<ExecutionSpace>, int> = 0>
auto reverse_copy(const ExecutionSpace& ex,
const ::Kokkos::View<DataType1, Properties1...>& source,
::Kokkos::View<DataType2, Properties2...>& dest) {
const ::Kokkos::View<DataType2, Properties2...>& dest) {
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(source);
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(dest);
@ -65,7 +65,7 @@ template <
std::enable_if_t<::Kokkos::is_execution_space_v<ExecutionSpace>, int> = 0>
auto reverse_copy(const std::string& label, const ExecutionSpace& ex,
const ::Kokkos::View<DataType1, Properties1...>& source,
::Kokkos::View<DataType2, Properties2...>& dest) {
const ::Kokkos::View<DataType2, Properties2...>& dest) {
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(source);
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(dest);
@ -94,7 +94,7 @@ template <typename TeamHandleType, typename DataType1, typename... Properties1,
KOKKOS_FUNCTION auto reverse_copy(
const TeamHandleType& teamHandle,
const ::Kokkos::View<DataType1, Properties1...>& source,
::Kokkos::View<DataType2, Properties2...>& dest) {
const ::Kokkos::View<DataType2, Properties2...>& dest) {
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(source);
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(dest);

41
lib/kokkos/algorithms/src/std_algorithms/Kokkos_Swap.hpp
View File

@ -1,41 +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_STD_ALGORITHMS_SWAP_HPP
#define KOKKOS_STD_ALGORITHMS_SWAP_HPP
#include <Kokkos_Core.hpp>
namespace Kokkos {
namespace Experimental {
// swap
template <class T>
KOKKOS_INLINE_FUNCTION void swap(T& a, T& b) noexcept {
static_assert(
std::is_move_assignable<T>::value && std::is_move_constructible<T>::value,
"Kokkos::Experimental::swap arguments must be move assignable "
"and move constructible");
T tmp = std::move(a);
a = std::move(b);
b = std::move(tmp);
}
} // namespace Experimental
} // namespace Kokkos
#endif

6
lib/kokkos/algorithms/src/std_algorithms/Kokkos_SwapRanges.hpp
View File

@ -40,7 +40,7 @@ template <typename ExecutionSpace, typename DataType1, typename... Properties1,
std::enable_if_t<is_execution_space_v<ExecutionSpace>, int> = 0>
auto swap_ranges(const ExecutionSpace& ex,
const ::Kokkos::View<DataType1, Properties1...>& source,
::Kokkos::View<DataType2, Properties2...>& dest) {
const ::Kokkos::View<DataType2, Properties2...>& dest) {
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(source);
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(dest);
@ -64,7 +64,7 @@ template <typename ExecutionSpace, typename DataType1, typename... Properties1,
std::enable_if_t<is_execution_space_v<ExecutionSpace>, int> = 0>
auto swap_ranges(const std::string& label, const ExecutionSpace& ex,
const ::Kokkos::View<DataType1, Properties1...>& source,
::Kokkos::View<DataType2, Properties2...>& dest) {
const ::Kokkos::View<DataType2, Properties2...>& dest) {
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(source);
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(dest);
@ -94,7 +94,7 @@ template <typename TeamHandleType, typename DataType1, typename... Properties1,
KOKKOS_FUNCTION auto swap_ranges(
const TeamHandleType& teamHandle,
const ::Kokkos::View<DataType1, Properties1...>& source,
::Kokkos::View<DataType2, Properties2...>& dest) {
const ::Kokkos::View<DataType2, Properties2...>& dest) {
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(source);
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(dest);

13
lib/kokkos/algorithms/src/std_algorithms/Kokkos_Transform.hpp
View File

@ -58,7 +58,7 @@ template <typename ExecutionSpace, typename DataType1, typename... Properties1,
std::enable_if_t<is_execution_space_v<ExecutionSpace>, int> = 0>
auto transform(const ExecutionSpace& ex,
const ::Kokkos::View<DataType1, Properties1...>& source,
::Kokkos::View<DataType2, Properties2...>& dest,
const ::Kokkos::View<DataType2, Properties2...>& dest,
UnaryOperation unary_op) {
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(source);
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(dest);
@ -73,7 +73,7 @@ template <typename ExecutionSpace, typename DataType1, typename... Properties1,
std::enable_if_t<is_execution_space_v<ExecutionSpace>, int> = 0>
auto transform(const std::string& label, const ExecutionSpace& ex,
const ::Kokkos::View<DataType1, Properties1...>& source,
::Kokkos::View<DataType2, Properties2...>& dest,
const ::Kokkos::View<DataType2, Properties2...>& dest,
UnaryOperation unary_op) {
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(source);
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(dest);
@ -119,7 +119,7 @@ template <typename ExecutionSpace, typename DataType1, typename... Properties1,
auto transform(const ExecutionSpace& ex,
const ::Kokkos::View<DataType1, Properties1...>& source1,
const ::Kokkos::View<DataType2, Properties2...>& source2,
::Kokkos::View<DataType3, Properties3...>& dest,
const ::Kokkos::View<DataType3, Properties3...>& dest,
BinaryOperation binary_op) {
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(source1);
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(source2);
@ -137,7 +137,7 @@ template <typename ExecutionSpace, typename DataType1, typename... Properties1,
auto transform(const std::string& label, const ExecutionSpace& ex,
const ::Kokkos::View<DataType1, Properties1...>& source1,
const ::Kokkos::View<DataType2, Properties2...>& source2,
::Kokkos::View<DataType3, Properties3...>& dest,
const ::Kokkos::View<DataType3, Properties3...>& dest,
BinaryOperation binary_op) {
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(source1);
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(source2);
@ -174,7 +174,8 @@ template <typename TeamHandleType, typename DataType1, typename... Properties1,
KOKKOS_FUNCTION auto transform(
const TeamHandleType& teamHandle,
const ::Kokkos::View<DataType1, Properties1...>& source,
::Kokkos::View<DataType2, Properties2...>& dest, UnaryOperation unary_op) {
const ::Kokkos::View<DataType2, Properties2...>& dest,
UnaryOperation unary_op) {
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(source);
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(dest);
@ -207,7 +208,7 @@ KOKKOS_FUNCTION auto transform(
const TeamHandleType& teamHandle,
const ::Kokkos::View<DataType1, Properties1...>& source1,
const ::Kokkos::View<DataType2, Properties2...>& source2,
::Kokkos::View<DataType3, Properties3...>& dest,
const ::Kokkos::View<DataType3, Properties3...>& dest,
BinaryOperation binary_op) {
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(source1);
Impl::static_assert_is_admissible_to_kokkos_std_algorithms(source2);

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

@ -47,8 +47,9 @@ struct ExclusiveScanDefaultFunctorForKnownNeutralElement {
KOKKOS_FUNCTION
void operator()(const IndexType i, ValueType& update,
const bool final_pass) const {
const auto tmp = m_first_from[i];
if (final_pass) m_first_dest[i] = update + m_init_value;
update += m_first_from[i];
update += tmp;
}
};
@ -73,6 +74,7 @@ struct ExclusiveScanDefaultFunctorWithValueWrapper {
KOKKOS_FUNCTION
void operator()(const IndexType i, value_type& update,
const bool final_pass) const {
const auto tmp = value_type{m_first_from[i], false};
if (final_pass) {
if (i == 0) {
m_first_dest[i] = m_init_value;
@ -81,7 +83,6 @@ struct ExclusiveScanDefaultFunctorWithValueWrapper {
}
}
const auto tmp = value_type{m_first_from[i], false};
this->join(update, tmp);
}
@ -132,6 +133,7 @@ struct TransformExclusiveScanFunctorWithValueWrapper {
KOKKOS_FUNCTION
void operator()(const IndexType i, value_type& update,
const bool final_pass) const {
const auto tmp = value_type{m_unary_op(m_first_from[i]), false};
if (final_pass) {
if (i == 0) {
// for both ExclusiveScan and TransformExclusiveScan,
@ -142,7 +144,6 @@ struct TransformExclusiveScanFunctorWithValueWrapper {
}
}
const auto tmp = value_type{m_unary_op(m_first_from[i]), false};
this->join(update, tmp);
}
@ -190,6 +191,7 @@ struct TransformExclusiveScanFunctorWithoutValueWrapper {
KOKKOS_FUNCTION
void operator()(const IndexType i, ValueType& update,
const bool final_pass) const {
const auto tmp = ValueType{m_unary_op(m_first_from[i])};
if (final_pass) {
if (i == 0) {
// for both ExclusiveScan and TransformExclusiveScan,
@ -200,7 +202,6 @@ struct TransformExclusiveScanFunctorWithoutValueWrapper {
}
}
const auto tmp = ValueType{m_unary_op(m_first_from[i])};
this->join(update, tmp);
}

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

@ -46,15 +46,14 @@ struct StdRemoveIfStage1Functor {
void operator()(const IndexType i, IndexType& update,
const bool final_pass) const {
auto& myval = m_first_from[i];
if (final_pass) {
if (!m_must_remove(myval)) {
if (!m_must_remove(myval)) {
if (final_pass) {
// calling move here is ok because we are inside final pass
// we are calling move assign as specified by the std
m_first_dest[update] = std::move(myval);
}
}
if (!m_must_remove(myval)) {
update += 1;
}
}
@ -108,7 +107,9 @@ IteratorType remove_if_exespace_impl(const std::string& label,
// create helper tmp view
using value_type = typename IteratorType::value_type;
using tmp_view_type = Kokkos::View<value_type*, ExecutionSpace>;
tmp_view_type tmp_view("std_remove_if_tmp_view", keep_count);
tmp_view_type tmp_view(Kokkos::view_alloc(Kokkos::WithoutInitializing, ex,
"std_remove_if_tmp_view"),
keep_count);
using tmp_readwrite_iterator_type = decltype(begin(tmp_view));
// in stage 1, *move* all elements to keep from original range to tmp

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

@ -21,7 +21,6 @@
#include "Kokkos_Constraints.hpp"
#include "Kokkos_HelperPredicates.hpp"
#include <std_algorithms/Kokkos_Distance.hpp>
#include <std_algorithms/Kokkos_Swap.hpp>
#include <string>
namespace Kokkos {
@ -39,7 +38,7 @@ struct StdReverseFunctor {
KOKKOS_FUNCTION
void operator()(index_type i) const {
::Kokkos::Experimental::swap(m_first[i], m_last[-i - 1]);
::Kokkos::kokkos_swap(m_first[i], m_last[-i - 1]);
}
KOKKOS_FUNCTION

5
lib/kokkos/algorithms/src/std_algorithms/impl/Kokkos_ShiftLeft.hpp
View File

@ -126,10 +126,11 @@ KOKKOS_FUNCTION IteratorType shift_left_team_impl(
// execution space impl because for this team impl we are
// within a parallel region, so for now we solve serially
const std::size_t numElementsToMove =
using difference_type = typename IteratorType::difference_type;
const difference_type numElementsToMove =
::Kokkos::Experimental::distance(first + n, last);
Kokkos::single(Kokkos::PerTeam(teamHandle), [=]() {
for (std::size_t i = 0; i < numElementsToMove; ++i) {
for (difference_type i = 0; i < numElementsToMove; ++i) {
first[i] = std::move(first[i + n]);
}
});

25
lib/kokkos/algorithms/src/std_algorithms/impl/Kokkos_ShiftRight.hpp
View File

@ -103,26 +103,6 @@ IteratorType shift_right_exespace_impl(
return first + n;
}
template <class Iterator>
struct StdShiftRightTeamSingleFunctor {
Iterator m_first;
Iterator m_last;
std::size_t m_shift;
KOKKOS_FUNCTION
void operator()() const {
// the impl function calling this functor guarantees that
// - m_shift is non-negative
// - m_first, m_last identify a valid range with m_last > m_first
// - m_shift is less than m_last - m_first
// so I can safely use std::size_t here
}
KOKKOS_FUNCTION
StdShiftRightTeamSingleFunctor(Iterator _first, Iterator _last, std::size_t n)
: m_first(std::move(_first)), m_last(std::move(_last)), m_shift(n) {}
};
template <class TeamHandleType, class IteratorType>
KOKKOS_FUNCTION IteratorType shift_right_team_impl(
const TeamHandleType& teamHandle, IteratorType first, IteratorType last,
@ -145,10 +125,11 @@ KOKKOS_FUNCTION IteratorType shift_right_team_impl(
// execution space impl because for this team impl we are
// within a parallel region, so for now we solve serially
const std::size_t numElementsToMove =
using difference_type = typename IteratorType::difference_type;
const difference_type numElementsToMove =
::Kokkos::Experimental::distance(first, last - n);
Kokkos::single(Kokkos::PerTeam(teamHandle), [=]() {
for (std::size_t i = 0; i < numElementsToMove; ++i) {
for (difference_type i = 0; i < numElementsToMove; ++i) {
last[-i - 1] = std::move(last[-n - i - 1]);
}
});

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

@ -21,7 +21,6 @@
#include "Kokkos_Constraints.hpp"
#include "Kokkos_HelperPredicates.hpp"
#include <std_algorithms/Kokkos_Distance.hpp>
#include <std_algorithms/Kokkos_Swap.hpp>
#include <string>
namespace Kokkos {
@ -36,7 +35,7 @@ struct StdSwapRangesFunctor {
KOKKOS_FUNCTION
void operator()(index_type i) const {
::Kokkos::Experimental::swap(m_first1[i], m_first2[i]);
::Kokkos::kokkos_swap(m_first1[i], m_first2[i]);
}
KOKKOS_FUNCTION

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

@ -105,7 +105,9 @@ IteratorType unique_exespace_impl(const std::string& label,
// using the same algorithm used for unique_copy but we now move things
using value_type = typename IteratorType::value_type;
using tmp_view_type = Kokkos::View<value_type*, ExecutionSpace>;
tmp_view_type tmp_view("std_unique_tmp_view", num_elements_to_explore);
tmp_view_type tmp_view(Kokkos::view_alloc(ex, Kokkos::WithoutInitializing,
"std_unique_tmp_view"),
num_elements_to_explore);
// scan extent is: num_elements_to_explore - 1
// for same reason as the one explained in unique_copy

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

@ -25,6 +25,7 @@ foreach(Tag Threads;Serial;OpenMP;Cuda;HPX;HIP;SYCL;OpenMPTarget)
set(ALGO_SORT_SOURCES)
foreach(SOURCE_Input
TestSort
TestSortByKey
TestSortCustomComp
TestBinSortA
TestBinSortB
@ -57,35 +58,37 @@ foreach(Tag Threads;Serial;OpenMP;Cuda;HPX;HIP;SYCL;OpenMPTarget)
configure_file(${dir}/dummy.cpp ${file})
list(APPEND ALGO_RANDOM_SOURCES ${file})
endforeach()
endif()
endforeach()
# ------------------------------------------
# std set A
# ------------------------------------------
set(STDALGO_SOURCES_A)
foreach(Name
# ------------------------------------------
# std set A
# ------------------------------------------
set(STDALGO_SOURCES_A)
foreach(Name
StdReducers
StdAlgorithmsConstraints
RandomAccessIterator
)
list(APPEND STDALGO_SOURCES_A Test${Name}.cpp)
endforeach()
)
list(APPEND STDALGO_SOURCES_A Test${Name}.cpp)
endforeach()
# ------------------------------------------
# std set B
# ------------------------------------------
set(STDALGO_SOURCES_B)
foreach(Name
# ------------------------------------------
# std set B
# ------------------------------------------
set(STDALGO_SOURCES_B)
foreach(Name
StdAlgorithmsCommon
StdAlgorithmsMinMaxElementOps
)
list(APPEND STDALGO_SOURCES_B Test${Name}.cpp)
endforeach()
)
list(APPEND STDALGO_SOURCES_B Test${Name}.cpp)
endforeach()
# ------------------------------------------
# std set C
# ------------------------------------------
set(STDALGO_SOURCES_C)
foreach(Name
# ------------------------------------------
# std set C
# ------------------------------------------
set(STDALGO_SOURCES_C)
foreach(Name
StdAlgorithmsCommon
StdAlgorithmsLexicographicalCompare
StdAlgorithmsForEach
@ -100,15 +103,15 @@ foreach(Tag Threads;Serial;OpenMP;Cuda;HPX;HIP;SYCL;OpenMPTarget)
StdAlgorithmsSearch_n
StdAlgorithmsMismatch
StdAlgorithmsMoveBackward
)
list(APPEND STDALGO_SOURCES_C Test${Name}.cpp)
endforeach()
)
list(APPEND STDALGO_SOURCES_C Test${Name}.cpp)
endforeach()
# ------------------------------------------
# std set D
# ------------------------------------------
set(STDALGO_SOURCES_D)
foreach(Name
# ------------------------------------------
# std set D
# ------------------------------------------
set(STDALGO_SOURCES_D)
foreach(Name
StdAlgorithmsCommon
StdAlgorithmsModOps
StdAlgorithmsModSeqOps
@ -128,15 +131,15 @@ foreach(Tag Threads;Serial;OpenMP;Cuda;HPX;HIP;SYCL;OpenMPTarget)
StdAlgorithmsReverse
StdAlgorithmsShiftLeft
StdAlgorithmsShiftRight
)
list(APPEND STDALGO_SOURCES_D Test${Name}.cpp)
endforeach()
)
list(APPEND STDALGO_SOURCES_D Test${Name}.cpp)
endforeach()
# ------------------------------------------
# std set E
# ------------------------------------------
set(STDALGO_SOURCES_E)
foreach(Name
# ------------------------------------------
# std set E
# ------------------------------------------
set(STDALGO_SOURCES_E)
foreach(Name
StdAlgorithmsCommon
StdAlgorithmsIsSorted
StdAlgorithmsIsSortedUntil
@ -149,83 +152,83 @@ foreach(Tag Threads;Serial;OpenMP;Cuda;HPX;HIP;SYCL;OpenMPTarget)
StdAlgorithmsTransformUnaryOp
StdAlgorithmsTransformExclusiveScan
StdAlgorithmsTransformInclusiveScan
)
list(APPEND STDALGO_SOURCES_E Test${Name}.cpp)
endforeach()
)
list(APPEND STDALGO_SOURCES_E Test${Name}.cpp)
endforeach()
# ------------------------------------------
# std team Q
# ------------------------------------------
set(STDALGO_TEAM_SOURCES_Q)
foreach(Name
# ------------------------------------------
# std team Q
# ------------------------------------------
set(STDALGO_TEAM_SOURCES_Q)
foreach(Name
StdAlgorithmsCommon
StdAlgorithmsTeamInclusiveScan
StdAlgorithmsTeamTransformInclusiveScan
)
list(APPEND STDALGO_TEAM_SOURCES_Q Test${Name}.cpp)
endforeach()
)
list(APPEND STDALGO_TEAM_SOURCES_Q Test${Name}.cpp)
endforeach()
# ------------------------------------------
# std team P
# ------------------------------------------
set(STDALGO_TEAM_SOURCES_P)
foreach(Name
# ------------------------------------------
# std team P
# ------------------------------------------
set(STDALGO_TEAM_SOURCES_P)
foreach(Name
StdAlgorithmsCommon
StdAlgorithmsTeamExclusiveScan
StdAlgorithmsTeamTransformExclusiveScan
)
list(APPEND STDALGO_TEAM_SOURCES_P Test${Name}.cpp)
endforeach()
)
list(APPEND STDALGO_TEAM_SOURCES_P Test${Name}.cpp)
endforeach()
# ------------------------------------------
# std team M
# ------------------------------------------
set(STDALGO_TEAM_SOURCES_M)
foreach(Name
# ------------------------------------------
# std team M
# ------------------------------------------
set(STDALGO_TEAM_SOURCES_M)
foreach(Name
StdAlgorithmsCommon
StdAlgorithmsTeamTransformUnaryOp
StdAlgorithmsTeamTransformBinaryOp
StdAlgorithmsTeamGenerate
StdAlgorithmsTeamGenerate_n
StdAlgorithmsTeamSwapRanges
)
list(APPEND STDALGO_TEAM_SOURCES_M Test${Name}.cpp)
endforeach()
)
list(APPEND STDALGO_TEAM_SOURCES_M Test${Name}.cpp)
endforeach()
# ------------------------------------------
# std team L
# ------------------------------------------
set(STDALGO_TEAM_SOURCES_L)
foreach(Name
# ------------------------------------------
# std team L
# ------------------------------------------
set(STDALGO_TEAM_SOURCES_L)
foreach(Name
StdAlgorithmsCommon
StdAlgorithmsTeamIsSorted
StdAlgorithmsTeamIsSortedUntil
StdAlgorithmsTeamIsPartitioned
StdAlgorithmsTeamPartitionCopy
StdAlgorithmsTeamPartitionPoint
)
list(APPEND STDALGO_TEAM_SOURCES_L Test${Name}.cpp)
endforeach()
)
list(APPEND STDALGO_TEAM_SOURCES_L Test${Name}.cpp)
endforeach()
# ------------------------------------------
# std team I
# ------------------------------------------
set(STDALGO_TEAM_SOURCES_I)
foreach(Name
# ------------------------------------------
# std team I
# ------------------------------------------
set(STDALGO_TEAM_SOURCES_I)
foreach(Name
StdAlgorithmsCommon
StdAlgorithmsTeamUnique
StdAlgorithmsTeamAdjacentDifference
StdAlgorithmsTeamReduce
StdAlgorithmsTeamTransformReduce
)
list(APPEND STDALGO_TEAM_SOURCES_I Test${Name}.cpp)
endforeach()
)
list(APPEND STDALGO_TEAM_SOURCES_I Test${Name}.cpp)
endforeach()
# ------------------------------------------
# std team H
# ------------------------------------------
set(STDALGO_TEAM_SOURCES_H)
foreach(Name
# ------------------------------------------
# std team H
# ------------------------------------------
set(STDALGO_TEAM_SOURCES_H)
foreach(Name
StdAlgorithmsCommon
StdAlgorithmsTeamCopy
StdAlgorithmsTeamCopy_n
@ -236,43 +239,43 @@ foreach(Tag Threads;Serial;OpenMP;Cuda;HPX;HIP;SYCL;OpenMPTarget)
StdAlgorithmsTeamRemoveIf
StdAlgorithmsTeamRemoveCopy
StdAlgorithmsTeamRemoveCopyIf
)
list(APPEND STDALGO_TEAM_SOURCES_H Test${Name}.cpp)
endforeach()
)
list(APPEND STDALGO_TEAM_SOURCES_H Test${Name}.cpp)
endforeach()
# ------------------------------------------
# std team G
# ------------------------------------------
set(STDALGO_TEAM_SOURCES_G)
foreach(Name
# ------------------------------------------
# std team G
# ------------------------------------------
set(STDALGO_TEAM_SOURCES_G)
foreach(Name
StdAlgorithmsCommon
StdAlgorithmsTeamMove
StdAlgorithmsTeamMoveBackward
StdAlgorithmsTeamShiftLeft
StdAlgorithmsTeamShiftRight
)
list(APPEND STDALGO_TEAM_SOURCES_G Test${Name}.cpp)
endforeach()
)
list(APPEND STDALGO_TEAM_SOURCES_G Test${Name}.cpp)
endforeach()
# ------------------------------------------
# std team F
# ------------------------------------------
set(STDALGO_TEAM_SOURCES_F)
foreach(Name
# ------------------------------------------
# std team F
# ------------------------------------------
set(STDALGO_TEAM_SOURCES_F)
foreach(Name
StdAlgorithmsCommon
StdAlgorithmsTeamReverse
StdAlgorithmsTeamReverseCopy
StdAlgorithmsTeamRotate
StdAlgorithmsTeamRotateCopy
)
list(APPEND STDALGO_TEAM_SOURCES_F Test${Name}.cpp)
endforeach()
)
list(APPEND STDALGO_TEAM_SOURCES_F Test${Name}.cpp)
endforeach()
# ------------------------------------------
# std team E
# ------------------------------------------
set(STDALGO_TEAM_SOURCES_E)
foreach(Name
# ------------------------------------------
# std team E
# ------------------------------------------
set(STDALGO_TEAM_SOURCES_E)
foreach(Name
StdAlgorithmsCommon
StdAlgorithmsTeamFill
StdAlgorithmsTeamFill_n
@ -280,28 +283,28 @@ foreach(Tag Threads;Serial;OpenMP;Cuda;HPX;HIP;SYCL;OpenMPTarget)
StdAlgorithmsTeamReplaceIf
StdAlgorithmsTeamReplaceCopy
StdAlgorithmsTeamReplaceCopyIf
)
list(APPEND STDALGO_TEAM_SOURCES_E Test${Name}.cpp)
endforeach()
)
list(APPEND STDALGO_TEAM_SOURCES_E Test${Name}.cpp)
endforeach()
# ------------------------------------------
# std team D
# ------------------------------------------
set(STDALGO_TEAM_SOURCES_D)
foreach(Name
# ------------------------------------------
# std team D
# ------------------------------------------
set(STDALGO_TEAM_SOURCES_D)
foreach(Name
StdAlgorithmsCommon
StdAlgorithmsTeamMinElement
StdAlgorithmsTeamMaxElement
StdAlgorithmsTeamMinMaxElement
)
list(APPEND STDALGO_TEAM_SOURCES_D Test${Name}.cpp)
endforeach()
)
list(APPEND STDALGO_TEAM_SOURCES_D Test${Name}.cpp)
endforeach()
# ------------------------------------------
# std team C
# ------------------------------------------
set(STDALGO_TEAM_SOURCES_C)
foreach(Name
# ------------------------------------------
# std team C
# ------------------------------------------
set(STDALGO_TEAM_SOURCES_C)
foreach(Name
StdAlgorithmsCommon
StdAlgorithmsTeamFind
StdAlgorithmsTeamFindIf
@ -310,29 +313,29 @@ foreach(Tag Threads;Serial;OpenMP;Cuda;HPX;HIP;SYCL;OpenMPTarget)
StdAlgorithmsTeamAnyOf
StdAlgorithmsTeamNoneOf
StdAlgorithmsTeamSearchN
)
list(APPEND STDALGO_TEAM_SOURCES_C Test${Name}.cpp)
endforeach()
)
list(APPEND STDALGO_TEAM_SOURCES_C Test${Name}.cpp)
endforeach()
# ------------------------------------------
# std team B
# ------------------------------------------
set(STDALGO_TEAM_SOURCES_B)
foreach(Name
# ------------------------------------------
# std team B
# ------------------------------------------
set(STDALGO_TEAM_SOURCES_B)
foreach(Name
StdAlgorithmsCommon
StdAlgorithmsTeamEqual
StdAlgorithmsTeamSearch
StdAlgorithmsTeamFindEnd
StdAlgorithmsTeamFindFirstOf
)
list(APPEND STDALGO_TEAM_SOURCES_B Test${Name}.cpp)
endforeach()
)
list(APPEND STDALGO_TEAM_SOURCES_B Test${Name}.cpp)
endforeach()
# ------------------------------------------
# std team A
# ------------------------------------------
set(STDALGO_TEAM_SOURCES_A)
foreach(Name
# ------------------------------------------
# std team A
# ------------------------------------------
set(STDALGO_TEAM_SOURCES_A)
foreach(Name
StdAlgorithmsCommon
StdAlgorithmsTeamAdjacentFind
StdAlgorithmsTeamCount
@ -341,11 +344,8 @@ foreach(Tag Threads;Serial;OpenMP;Cuda;HPX;HIP;SYCL;OpenMPTarget)
StdAlgorithmsTeamForEachN
StdAlgorithmsTeamLexicographicalCompare
StdAlgorithmsTeamMismatch
)
list(APPEND STDALGO_TEAM_SOURCES_A Test${Name}.cpp)
endforeach()
endif()
)
list(APPEND STDALGO_TEAM_SOURCES_A Test${Name}.cpp)
endforeach()
# FIXME_OPENMPTARGET - remove sort test as it leads to ICE with clang/16 and above at compile time.

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

@ -27,13 +27,13 @@ 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 <TestRandom.hpp>" >> Test$(device).cpp); \
$(shell echo "\#include <TestSort.hpp>" >> Test$(device).cpp); \
$(shell echo "\#include <TestBinSortA.hpp>" >> Test$(device).cpp); \
$(shell echo "\#include <TestBinSortB.hpp>" >> Test$(device).cpp); \
$(shell echo "\#include <TestNestedSort.hpp>" >> Test$(device).cpp); \
$(shell echo "\#include <TestSortCustomComp.hpp>" >> Test$(device).cpp); \
$(shell echo "$(H)include <Test"${device}"_Category.hpp>" > Test$(device).cpp); \
$(shell echo "$(H)include <TestRandom.hpp>" >> Test$(device).cpp); \
$(shell echo "$(H)include <TestSort.hpp>" >> Test$(device).cpp); \
$(shell echo "$(H)include <TestBinSortA.hpp>" >> Test$(device).cpp); \
$(shell echo "$(H)include <TestBinSortB.hpp>" >> Test$(device).cpp); \
$(shell echo "$(H)include <TestNestedSort.hpp>" >> Test$(device).cpp); \
$(shell echo "$(H)include <TestSortCustomComp.hpp>" >> Test$(device).cpp); \
) \
)

241
lib/kokkos/algorithms/unit_tests/TestSortByKey.hpp Normal file
View File

@ -0,0 +1,241 @@
//@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_SORT_BY_KEY_HPP
#define KOKKOS_ALGORITHMS_UNITTESTS_TEST_SORT_BY_KEY_HPP
#include <gtest/gtest.h>
#include <Kokkos_Core.hpp>
#include <Kokkos_Random.hpp>
#include <Kokkos_Sort.hpp>
#include <utility> // pair
namespace Test {
namespace SortImpl {
struct Less {
template <class ValueType>
KOKKOS_INLINE_FUNCTION bool operator()(const ValueType &lhs,
const ValueType &rhs) const {
return lhs < rhs;
}
};
struct Greater {
template <class ValueType>
KOKKOS_INLINE_FUNCTION bool operator()(const ValueType &lhs,
const ValueType &rhs) const {
return lhs > rhs;
}
};
template <class ExecutionSpace, class Keys, class Permute,
class Comparator = Less>
struct is_sorted_by_key_struct {
Keys keys;
Keys keys_orig;
Permute permute;
Comparator comparator;
is_sorted_by_key_struct(Keys keys_, Keys keys_orig_, Permute permute_,
Comparator comparator_ = Comparator{})
: keys(keys_),
keys_orig(keys_orig_),
permute(permute_),
comparator(comparator_) {}
KOKKOS_INLINE_FUNCTION
void operator()(int i, unsigned int &count) const {
if (i < keys.extent_int(0) - 1 && comparator(keys(i + 1), keys(i))) ++count;
if (keys(i) != keys_orig(permute(i))) ++count;
}
};
template <typename ExecutionSpace, typename ViewType>
void iota(ExecutionSpace const &space, ViewType const &v,
typename ViewType::value_type value = 0) {
using ValueType = typename ViewType::value_type;
Kokkos::parallel_for(
"ArborX::Algorithms::iota",
Kokkos::RangePolicy<ExecutionSpace>(space, 0, v.extent(0)),
KOKKOS_LAMBDA(int i) { v(i) = value + (ValueType)i; });
}
} // namespace SortImpl
TEST(TEST_CATEGORY, SortByKeyEmptyView) {
using ExecutionSpace = TEST_EXECSPACE;
// does not matter if we use int or something else
Kokkos::View<int *, ExecutionSpace> keys("keys", 0);
Kokkos::View<float *, ExecutionSpace> values("values", 0);
ASSERT_NO_THROW(
Kokkos::Experimental::sort_by_key(ExecutionSpace(), keys, values));
}
TEST(TEST_CATEGORY, SortByKey) {
using ExecutionSpace = TEST_EXECSPACE;
using MemorySpace = typename ExecutionSpace::memory_space;
ExecutionSpace space{};
for (auto keys_vector : {std::vector<int>{36, 19, 25, 17, 3, 7, 1, 2, 9},
std::vector<int>{36, 19, 25, 17, 3, 9, 1, 2, 7},
std::vector<int>{100, 19, 36, 17, 3, 25, 1, 2, 7},
std::vector<int>{15, 5, 11, 3, 4, 8}}) {
auto const n = keys_vector.size();
auto keys = Kokkos::create_mirror_view_and_copy(
MemorySpace{},
Kokkos::View<int *, Kokkos::HostSpace, Kokkos::MemoryUnmanaged>(
keys_vector.data(), n));
auto keys_orig = Kokkos::create_mirror(space, keys);
Kokkos::deep_copy(space, keys_orig, keys);
Kokkos::View<int *, ExecutionSpace> permute("permute", n);
SortImpl::iota(space, permute);
Kokkos::Experimental::sort_by_key(space, keys, permute);
unsigned int sort_fails = 0;
Kokkos::parallel_reduce(
Kokkos::RangePolicy<ExecutionSpace>(space, 0, n),
SortImpl::is_sorted_by_key_struct<ExecutionSpace, decltype(keys),
decltype(permute)>(keys, keys_orig,
permute),
sort_fails);
ASSERT_EQ(sort_fails, 0u);
}
}
TEST(TEST_CATEGORY, SortByKeyWithComparator) {
using ExecutionSpace = TEST_EXECSPACE;
using MemorySpace = typename ExecutionSpace::memory_space;
ExecutionSpace space{};
SortImpl::Greater comparator;
for (auto keys_vector : {std::vector<int>{36, 19, 25, 17, 3, 7, 1, 2, 9},
std::vector<int>{36, 19, 25, 17, 3, 9, 1, 2, 7},
std::vector<int>{100, 19, 36, 17, 3, 25, 1, 2, 7},
std::vector<int>{15, 5, 11, 3, 4, 8}}) {
auto const n = keys_vector.size();
auto keys = Kokkos::create_mirror_view_and_copy(
MemorySpace{},
Kokkos::View<int *, Kokkos::HostSpace, Kokkos::MemoryUnmanaged>(
keys_vector.data(), n));
auto keys_orig = Kokkos::create_mirror(space, keys);
Kokkos::deep_copy(space, keys_orig, keys);
Kokkos::View<int *, ExecutionSpace> permute("permute", n);
SortImpl::iota(space, permute);
Kokkos::Experimental::sort_by_key(space, keys, permute, comparator);
unsigned int sort_fails = 0;
Kokkos::parallel_reduce(
Kokkos::RangePolicy<ExecutionSpace>(space, 0, n),
SortImpl::is_sorted_by_key_struct<ExecutionSpace, decltype(keys),
decltype(permute), SortImpl::Greater>(
keys, keys_orig, permute, comparator),
sort_fails);
ASSERT_EQ(sort_fails, 0u);
}
}
TEST(TEST_CATEGORY, SortByKeyStaticExtents) {
using ExecutionSpace = TEST_EXECSPACE;
ExecutionSpace space{};
Kokkos::View<int[10], ExecutionSpace> keys("keys");
Kokkos::View<int[10], ExecutionSpace> values_static("values_static");
ASSERT_NO_THROW(
Kokkos::Experimental::sort_by_key(space, keys, values_static));
Kokkos::View<int *, ExecutionSpace> values_dynamic("values_dynamic", 10);
ASSERT_NO_THROW(
Kokkos::Experimental::sort_by_key(space, keys, values_dynamic));
}
template <typename ExecutionSpace, typename Keys, typename Values>
void buildViewsForStrided(ExecutionSpace const &space, int n, Keys &keys,
Values &values) {
Kokkos::parallel_for(
"create_data",
Kokkos::MDRangePolicy<Kokkos::Rank<3>, ExecutionSpace>(space, {0, 0, 0},
{n, n, n}),
KOKKOS_LAMBDA(int i, int j, int k) {
keys(i, j, k) = n - i;
values(i, j, k) = j;
});
}
TEST(TEST_CATEGORY, SortByKeyWithStrides) {
using ExecutionSpace = TEST_EXECSPACE;
ExecutionSpace space{};
auto const n = 10;
Kokkos::View<int ***, ExecutionSpace> keys("keys", n, n, n);
Kokkos::View<int ***, ExecutionSpace> values("values", n, n, n);
buildViewsForStrided(space, n, keys, values);
auto keys_sub = Kokkos::subview(keys, Kokkos::ALL(), 1, 2);
auto values_sub = Kokkos::subview(values, 4, Kokkos::ALL(), 6);
auto keys_orig = Kokkos::create_mirror(space, keys_sub);
Kokkos::deep_copy(space, keys_orig, keys_sub);
Kokkos::Experimental::sort_by_key(space, keys_sub, values_sub);
unsigned int sort_fails = 0;
Kokkos::parallel_reduce(
Kokkos::RangePolicy<ExecutionSpace>(space, 0, n),
SortImpl::is_sorted_by_key_struct<ExecutionSpace, decltype(keys_sub),
decltype(values_sub)>(
keys_sub, keys_orig, values_sub),
sort_fails);
ASSERT_EQ(sort_fails, 0u);
}
TEST(TEST_CATEGORY, SortByKeyKeysLargerThanValues) {
using ExecutionSpace = TEST_EXECSPACE;
// does not matter if we use int or something else
Kokkos::View<int *, ExecutionSpace> keys("keys", 3);
Kokkos::View<float *, ExecutionSpace> values("values", 1);
ASSERT_DEATH(
Kokkos::Experimental::sort_by_key(ExecutionSpace(), keys, values),
"values and keys extents must be the same");
ASSERT_DEATH(Kokkos::Experimental::sort_by_key(ExecutionSpace(), keys, values,
SortImpl::Greater{}),
"values and keys extents must be the same");
}
} // namespace Test
#endif

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

@ -239,16 +239,8 @@ KOKKOS_FUNCTION bool team_members_have_matching_result(
// set accum to 1 if a mismach is found
const bool mismatch = memberValue != target;
int accum = static_cast<int>(mismatch);
// FIXME_OPENMPTARGET: team API does not meet the TeamHandle concept and
// ignores the reducer passed
#if defined KOKKOS_ENABLE_OPENMPTARGET
Kokkos::Sum<int> dummyReducer(accum);
const auto result = teamHandle.team_reduce(accum, dummyReducer);
return (result == 0);
#else
teamHandle.team_reduce(Kokkos::Sum<int>(accum));
return (accum == 0);
#endif
}
template <class ValueType1, class ValueType2>

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

@ -16,6 +16,7 @@
#include <TestStdAlgorithmsCommon.hpp>
#include <utility>
#include <iomanip>
namespace Test {
namespace stdalgos {
@ -132,47 +133,6 @@ void my_host_exclusive_scan(it1 first, it1 last, it2 dest, ValType init,
}
}
template <class ViewType1, class ViewType2, class ValueType, class BinaryOp>
void verify_data(ViewType1 data_view, // contains data
ViewType2 test_view, // the view to test
ValueType init_value, BinaryOp bop) {
//! always careful because views might not be deep copyable
auto data_view_dc = create_deep_copyable_compatible_clone(data_view);
auto data_view_h =
create_mirror_view_and_copy(Kokkos::HostSpace(), data_view_dc);
using gold_view_value_type = typename ViewType2::value_type;
Kokkos::View<gold_view_value_type*, Kokkos::HostSpace> gold_h(
"goldh", data_view.extent(0));
my_host_exclusive_scan(KE::cbegin(data_view_h), KE::cend(data_view_h),
KE::begin(gold_h), init_value, bop);
auto test_view_dc = create_deep_copyable_compatible_clone(test_view);
auto test_view_h =
create_mirror_view_and_copy(Kokkos::HostSpace(), test_view_dc);
if (test_view_h.extent(0) > 0) {
for (std::size_t i = 0; i < test_view_h.extent(0); ++i) {
// std::cout << i << " " << std::setprecision(15) << data_view_h(i) << " "
// << 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) {
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)));
if (error > 1e-10) {
std::cout << i << " " << std::setprecision(15) << data_view_h(i)
<< " " << gold_h(i) << " " << test_view_h(i) << " "
<< std::abs(static_cast<double>(gold_h(i) - test_view_h(i)))
<< std::endl;
}
EXPECT_LT(error, 1e-10);
}
}
}
}
template <class ValueType>
struct MultiplyFunctor {
KOKKOS_INLINE_FUNCTION
@ -189,107 +149,153 @@ struct SumFunctor {
}
};
struct VerifyData {
template <class ViewType1, class ViewType2, class ValueType, class BinaryOp>
void operator()(ViewType1 data_view, // contains data
ViewType2 test_view, // the view to test
ValueType init_value, BinaryOp bop) {
//! always careful because views might not be deep copyable
auto data_view_dc = create_deep_copyable_compatible_clone(data_view);
auto data_view_h =
create_mirror_view_and_copy(Kokkos::HostSpace(), data_view_dc);
using gold_view_value_type = typename ViewType2::value_type;
Kokkos::View<gold_view_value_type*, Kokkos::HostSpace> gold_h(
"goldh", data_view.extent(0));
my_host_exclusive_scan(KE::cbegin(data_view_h), KE::cend(data_view_h),
KE::begin(gold_h), init_value, bop);
auto test_view_dc = create_deep_copyable_compatible_clone(test_view);
auto test_view_h =
create_mirror_view_and_copy(Kokkos::HostSpace(), test_view_dc);
if (test_view_h.extent(0) > 0) {
for (std::size_t i = 0; i < test_view_h.extent(0); ++i) {
if (std::is_same<gold_view_value_type, int>::value) {
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)));
ASSERT_LT(error, 1e-10) << i << " " << std::setprecision(15) << error
<< static_cast<double>(test_view_h(i)) << " "
<< static_cast<double>(gold_h(i));
}
}
}
}
template <class ViewType1, class ViewType2, class ValueType>
void operator()(ViewType1 data_view, // contains data
ViewType2 test_view, // the view to test
ValueType init_value) {
(*this)(data_view, test_view, init_value, SumFunctor<ValueType>());
}
};
std::string value_type_to_string(int) { return "int"; }
std::string value_type_to_string(double) { return "double"; }
template <class Tag, class ValueType, class InfoType>
void run_single_scenario_default_op(const InfoType& scenario_info,
ValueType init_value) {
using default_op = SumFunctor<ValueType>;
template <class Tag, class ValueType, class InfoType, class... OpOrEmpty>
void run_single_scenario(const InfoType& scenario_info, ValueType init_value,
OpOrEmpty... empty_or_op) {
const auto name = std::get<0>(scenario_info);
const std::size_t view_ext = std::get<1>(scenario_info);
// std::cout << "exclusive_scan default op: " << name << ", "
// << view_tag_to_string(Tag{}) << ", "
// << value_type_to_string(ValueType()) << ", "
// << "init = " << init_value << std::endl;
auto view_dest = create_view<ValueType>(Tag{}, view_ext, "exclusive_scan");
auto view_from = create_view<ValueType>(Tag{}, view_ext, "exclusive_scan");
fill_view(view_from, name);
// view_dest is filled with zeros before calling the algorithm everytime to
// ensure the algorithm does something meaningful
{
fill_zero(view_dest);
auto r = KE::exclusive_scan(exespace(), KE::cbegin(view_from),
KE::cend(view_from), KE::begin(view_dest),
init_value);
init_value, empty_or_op...);
ASSERT_EQ(r, KE::end(view_dest));
verify_data(view_from, view_dest, init_value, default_op());
VerifyData()(view_from, view_dest, init_value, empty_or_op...);
}
{
fill_zero(view_dest);
auto r = KE::exclusive_scan("label", exespace(), KE::cbegin(view_from),
KE::cend(view_from), KE::begin(view_dest),
init_value);
init_value, empty_or_op...);
ASSERT_EQ(r, KE::end(view_dest));
verify_data(view_from, view_dest, init_value, default_op());
VerifyData()(view_from, view_dest, init_value, empty_or_op...);
}
{
fill_zero(view_dest);
auto r = KE::exclusive_scan(exespace(), view_from, view_dest, init_value);
auto r = KE::exclusive_scan(exespace(), view_from, view_dest, init_value,
empty_or_op...);
ASSERT_EQ(r, KE::end(view_dest));
verify_data(view_from, view_dest, init_value, default_op());
VerifyData()(view_from, view_dest, init_value, empty_or_op...);
}
{
fill_zero(view_dest);
auto r = KE::exclusive_scan("label", exespace(), view_from, view_dest,
init_value);
init_value, empty_or_op...);
ASSERT_EQ(r, KE::end(view_dest));
verify_data(view_from, view_dest, init_value, default_op());
VerifyData()(view_from, view_dest, init_value, empty_or_op...);
}
Kokkos::fence();
}
template <class Tag, class ValueType, class InfoType, class BinaryOp>
void run_single_scenario_custom_op(const InfoType& scenario_info,
ValueType init_value, BinaryOp bop) {
template <class Tag, class ValueType, class InfoType, class... OpOrEmpty>
void run_single_scenario_inplace(const InfoType& scenario_info,
ValueType init_value,
OpOrEmpty... empty_or_op) {
const auto name = std::get<0>(scenario_info);
const std::size_t view_ext = std::get<1>(scenario_info);
// std::cout << "exclusive_scan custom op: " << name << ", "
// << view_tag_to_string(Tag{}) << ", "
// << value_type_to_string(ValueType()) << ", "
// << "init = " << init_value << std::endl;
auto view_dest = create_view<ValueType>(Tag{}, view_ext, "exclusive_scan");
auto view_from = create_view<ValueType>(Tag{}, view_ext, "exclusive_scan");
fill_view(view_from, name);
// since here we call the in-place operation, we need to use two views:
// view1: filled according to what the scenario asks for and is not modified
// view2: filled according to what the scenario asks for and used for the
// in-place op Therefore, after the op is done, view2 should contain the
// result of doing exclusive scan NOTE: view2 is filled below every time
// because the algorithm acts in place
auto view1 =
create_view<ValueType>(Tag{}, view_ext, "exclusive_scan_inplace_view1");
fill_view(view1, name);
auto view2 =
create_view<ValueType>(Tag{}, view_ext, "exclusive_scan_inplace_view2");
{
fill_zero(view_dest);
auto r = KE::exclusive_scan(exespace(), KE::cbegin(view_from),
KE::cend(view_from), KE::begin(view_dest),
init_value, bop);
ASSERT_EQ(r, KE::end(view_dest));
verify_data(view_from, view_dest, init_value, bop);
fill_view(view2, name);
auto r = KE::exclusive_scan(exespace(), KE::cbegin(view2), KE::cend(view2),
KE::begin(view2), init_value, empty_or_op...);
ASSERT_EQ(r, KE::end(view2));
VerifyData()(view1, view2, init_value, empty_or_op...);
}
{
fill_zero(view_dest);
auto r = KE::exclusive_scan("label", exespace(), KE::cbegin(view_from),
KE::cend(view_from), KE::begin(view_dest),
init_value, bop);
ASSERT_EQ(r, KE::end(view_dest));
verify_data(view_from, view_dest, init_value, bop);
fill_view(view2, name);
auto r = KE::exclusive_scan("label", exespace(), KE::cbegin(view2),
KE::cend(view2), KE::begin(view2), init_value,
empty_or_op...);
ASSERT_EQ(r, KE::end(view2));
VerifyData()(view1, view2, init_value, empty_or_op...);
}
{
fill_zero(view_dest);
auto r =
KE::exclusive_scan(exespace(), view_from, view_dest, init_value, bop);
ASSERT_EQ(r, KE::end(view_dest));
verify_data(view_from, view_dest, init_value, bop);
fill_view(view2, name);
auto r = KE::exclusive_scan(exespace(), view2, view2, init_value,
empty_or_op...);
ASSERT_EQ(r, KE::end(view2));
VerifyData()(view1, view2, init_value, empty_or_op...);
}
{
fill_zero(view_dest);
auto r = KE::exclusive_scan("label", exespace(), view_from, view_dest,
init_value, bop);
ASSERT_EQ(r, KE::end(view_dest));
verify_data(view_from, view_dest, init_value, bop);
fill_view(view2, name);
auto r = KE::exclusive_scan("label", exespace(), view2, view2, init_value,
empty_or_op...);
ASSERT_EQ(r, KE::end(view2));
VerifyData()(view1, view2, init_value, empty_or_op...);
}
Kokkos::fence();
@ -303,34 +309,39 @@ void run_exclusive_scan_all_scenarios() {
{"medium", 1103}, {"large", 10513}};
for (const auto& it : scenarios) {
run_single_scenario_default_op<Tag, ValueType>(it, ValueType{0});
run_single_scenario_default_op<Tag, ValueType>(it, ValueType{1});
run_single_scenario_default_op<Tag, ValueType>(it, ValueType{-2});
run_single_scenario_default_op<Tag, ValueType>(it, ValueType{3});
run_single_scenario<Tag, ValueType>(it, ValueType{0});
run_single_scenario<Tag, ValueType>(it, ValueType{1});
run_single_scenario<Tag, ValueType>(it, ValueType{-2});
run_single_scenario<Tag, ValueType>(it, ValueType{3});
run_single_scenario_inplace<Tag, ValueType>(it, ValueType{0});
run_single_scenario_inplace<Tag, ValueType>(it, ValueType{-2});
#if !defined KOKKOS_ENABLE_OPENMPTARGET
// custom multiply op is only run for small views otherwise it overflows
if (it.first == "small-a" || it.first == "small-b") {
using custom_bop_t = MultiplyFunctor<ValueType>;
run_single_scenario_custom_op<Tag, ValueType>(it, ValueType{0},
custom_bop_t());
run_single_scenario_custom_op<Tag, ValueType>(it, ValueType{1},
custom_bop_t());
run_single_scenario_custom_op<Tag, ValueType>(it, ValueType{-2},
custom_bop_t());
run_single_scenario_custom_op<Tag, ValueType>(it, ValueType{3},
custom_bop_t());
run_single_scenario<Tag, ValueType>(it, ValueType{0}, custom_bop_t());
run_single_scenario<Tag, ValueType>(it, ValueType{1}, custom_bop_t());
run_single_scenario<Tag, ValueType>(it, ValueType{-2}, custom_bop_t());
run_single_scenario<Tag, ValueType>(it, ValueType{3}, custom_bop_t());
run_single_scenario_inplace<Tag, ValueType>(it, ValueType{0},
custom_bop_t());
run_single_scenario_inplace<Tag, ValueType>(it, ValueType{-2},
custom_bop_t());
}
using custom_bop_t = SumFunctor<ValueType>;
run_single_scenario_custom_op<Tag, ValueType>(it, ValueType{0},
custom_bop_t());
run_single_scenario_custom_op<Tag, ValueType>(it, ValueType{1},
custom_bop_t());
run_single_scenario_custom_op<Tag, ValueType>(it, ValueType{-2},
custom_bop_t());
run_single_scenario_custom_op<Tag, ValueType>(it, ValueType{3},
custom_bop_t());
run_single_scenario<Tag, ValueType>(it, ValueType{0}, custom_bop_t());
run_single_scenario<Tag, ValueType>(it, ValueType{1}, custom_bop_t());
run_single_scenario<Tag, ValueType>(it, ValueType{-2}, custom_bop_t());
run_single_scenario<Tag, ValueType>(it, ValueType{3}, custom_bop_t());
run_single_scenario_inplace<Tag, ValueType>(it, ValueType{0},
custom_bop_t());
run_single_scenario_inplace<Tag, ValueType>(it, ValueType{-2},
custom_bop_t());
#endif
}
}

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

@ -16,6 +16,7 @@
#include <TestStdAlgorithmsCommon.hpp>
#include <utility>
#include <iomanip>
namespace Test {
namespace stdalgos {
@ -143,51 +144,6 @@ void my_host_inclusive_scan(it1 first, it1 last, it2 dest, BinOp bop,
}
}
template <class ViewType1, class ViewType2, class BinaryOp, class... Args>
void verify_data(ViewType1 data_view, // contains data
ViewType2 test_view, // the view to test
BinaryOp bop, Args... args /* copy on purpose */) {
//! always careful because views might not be deep copyable
auto data_view_dc = create_deep_copyable_compatible_clone(data_view);
auto data_view_h =
create_mirror_view_and_copy(Kokkos::HostSpace(), data_view_dc);
using gold_view_value_type = typename ViewType2::value_type;
Kokkos::View<gold_view_value_type*, Kokkos::HostSpace> gold_h(
"goldh", data_view.extent(0));
my_host_inclusive_scan(KE::cbegin(data_view_h), KE::cend(data_view_h),
KE::begin(gold_h), bop, args...);
auto test_view_dc = create_deep_copyable_compatible_clone(test_view);
auto test_view_h =
create_mirror_view_and_copy(Kokkos::HostSpace(), test_view_dc);
const auto ext = test_view_h.extent(0);
if (ext > 0) {
for (std::size_t i = 0; i < ext; ++i) {
// std::cout << i << " " << std::setprecision(15) << data_view_h(i) << " "
// << 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) {
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)));
if (error > 1e-10) {
std::cout << i << " " << std::setprecision(15) << data_view_h(i)
<< " " << gold_h(i) << " " << test_view_h(i) << " "
<< std::abs(static_cast<double>(gold_h(i) - test_view_h(i)))
<< std::endl;
}
EXPECT_LT(error, 1e-10);
}
}
// std::cout << " last el: " << test_view_h(ext-1) << std::endl;
}
}
template <class ValueType>
struct MultiplyFunctor {
KOKKOS_INLINE_FUNCTION
@ -204,107 +160,151 @@ struct SumFunctor {
}
};
struct VerifyData {
template <class ViewType1, class ViewType2, class BinaryOp, class... Args>
void operator()(ViewType1 data_view, // contains data
ViewType2 test_view, // the view to test
BinaryOp bop, Args... args /* copy on purpose */) {
//! always careful because views might not be deep copyable
auto data_view_dc = create_deep_copyable_compatible_clone(data_view);
auto data_view_h =
create_mirror_view_and_copy(Kokkos::HostSpace(), data_view_dc);
using gold_view_value_type = typename ViewType2::value_type;
Kokkos::View<gold_view_value_type*, Kokkos::HostSpace> gold_h(
"goldh", data_view.extent(0));
my_host_inclusive_scan(KE::cbegin(data_view_h), KE::cend(data_view_h),
KE::begin(gold_h), bop, args...);
auto test_view_dc = create_deep_copyable_compatible_clone(test_view);
auto test_view_h =
create_mirror_view_and_copy(Kokkos::HostSpace(), test_view_dc);
const auto ext = test_view_h.extent(0);
if (ext > 0) {
for (std::size_t i = 0; i < ext; ++i) {
if (std::is_same<gold_view_value_type, int>::value) {
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)));
ASSERT_LT(error, 1e-10) << i << " " << std::setprecision(15) << error
<< static_cast<double>(test_view_h(i)) << " "
<< static_cast<double>(gold_h(i));
}
}
}
}
template <class ViewType1, class ViewType2>
void operator()(ViewType1 data_view, // contains data
ViewType2 test_view) // the view to test
{
using value_type = typename ViewType1::non_const_value_type;
(*this)(data_view, test_view, SumFunctor<value_type>());
}
};
std::string value_type_to_string(int) { return "int"; }
std::string value_type_to_string(double) { return "double"; }
template <class Tag, class ValueType, class InfoType>
void run_single_scenario_default_op(const InfoType& scenario_info) {
using default_op = SumFunctor<ValueType>;
template <class Tag, class ValueType, class InfoType, class... Args>
void run_single_scenario(const InfoType& scenario_info,
Args... args /* copy on purpose */) {
const auto name = std::get<0>(scenario_info);
const std::size_t view_ext = std::get<1>(scenario_info);
// std::cout << "inclusive_scan default op: " << name << ", "
// << view_tag_to_string(Tag{}) << ", "
// << value_type_to_string(ValueType()) << std::endl;
auto view_dest = create_view<ValueType>(Tag{}, view_ext, "inclusive_scan");
auto view_from = create_view<ValueType>(Tag{}, view_ext, "inclusive_scan");
fill_view(view_from, name);
// view_dest is filled with zeros before calling the algorithm everytime to
// ensure the algorithm does something meaningful
{
fill_zero(view_dest);
auto r = KE::inclusive_scan(exespace(), KE::cbegin(view_from),
KE::cend(view_from), KE::begin(view_dest));
auto r =
KE::inclusive_scan(exespace(), KE::cbegin(view_from),
KE::cend(view_from), KE::begin(view_dest), args...);
ASSERT_EQ(r, KE::end(view_dest));
verify_data(view_from, view_dest, default_op());
VerifyData()(view_from, view_dest, args...);
}
{
fill_zero(view_dest);
auto r = KE::inclusive_scan("label", exespace(), KE::cbegin(view_from),
KE::cend(view_from), KE::begin(view_dest));
auto r =
KE::inclusive_scan("label", exespace(), KE::cbegin(view_from),
KE::cend(view_from), KE::begin(view_dest), args...);
ASSERT_EQ(r, KE::end(view_dest));
verify_data(view_from, view_dest, default_op());
VerifyData()(view_from, view_dest, args...);
}
{
fill_zero(view_dest);
auto r = KE::inclusive_scan(exespace(), view_from, view_dest);
auto r = KE::inclusive_scan(exespace(), view_from, view_dest, args...);
ASSERT_EQ(r, KE::end(view_dest));
verify_data(view_from, view_dest, default_op());
VerifyData()(view_from, view_dest, args...);
}
{
fill_zero(view_dest);
auto r = KE::inclusive_scan("label", exespace(), view_from, view_dest);
auto r =
KE::inclusive_scan("label", exespace(), view_from, view_dest, args...);
ASSERT_EQ(r, KE::end(view_dest));
verify_data(view_from, view_dest, default_op());
VerifyData()(view_from, view_dest, args...);
}
Kokkos::fence();
}
template <class Tag, class ValueType, class InfoType, class BinaryOp,
class... Args>
void run_single_scenario_custom_op(const InfoType& scenario_info, BinaryOp bop,
Args... args /* copy on purpose */) {
template <class Tag, class ValueType, class InfoType, class... Args>
void run_single_scenario_inplace(const InfoType& scenario_info,
Args... args /* copy on purpose */) {
const auto name = std::get<0>(scenario_info);
const std::size_t view_ext = std::get<1>(scenario_info);
// if (1 == sizeof...(Args)) {
// std::cout << "inclusive_scan custom op and init value: " << name << ", "
// << view_tag_to_string(Tag{}) << ", "
// << value_type_to_string(ValueType()) << ", " << std::endl;
// } else {
// std::cout << "inclusive_scan custom op: " << name << ", "
// << view_tag_to_string(Tag{}) << ", "
// << value_type_to_string(ValueType()) << ", " << std::endl;
// }
// since here we call the in-place operation, we need to use two views:
// view1: filled according to what the scenario asks for and is not modified
// view2: filled according to what the scenario asks for and used for the
// in-place op Therefore, after the op is done, view_2 should contain the
// result of doing exclusive scan NOTE: view2 is filled below every time
// because the algorithm acts in place
auto view_dest = create_view<ValueType>(Tag{}, view_ext, "inclusive_scan");
auto view_from = create_view<ValueType>(Tag{}, view_ext, "inclusive_scan");
fill_view(view_from, name);
auto view1 =
create_view<ValueType>(Tag{}, view_ext, "inclusive_scan_inplace_view1");
fill_view(view1, name);
auto view2 =
create_view<ValueType>(Tag{}, view_ext, "inclusive_scan_inplace_view2");
{
fill_zero(view_dest);
auto r = KE::inclusive_scan(exespace(), KE::cbegin(view_from),
KE::cend(view_from), KE::begin(view_dest), bop,
args...);
ASSERT_EQ(r, KE::end(view_dest));
verify_data(view_from, view_dest, bop, args...);
fill_view(view2, name);
auto r = KE::inclusive_scan(exespace(), KE::cbegin(view2), KE::cend(view2),
KE::begin(view2), args...);
ASSERT_EQ(r, KE::end(view2));
VerifyData()(view1, view2, args...);
}
{
fill_zero(view_dest);
auto r = KE::inclusive_scan("label", exespace(), KE::cbegin(view_from),
KE::cend(view_from), KE::begin(view_dest), bop,
args...);
ASSERT_EQ(r, KE::end(view_dest));
verify_data(view_from, view_dest, bop, args...);
fill_view(view2, name);
auto r = KE::inclusive_scan("label", exespace(), KE::cbegin(view2),
KE::cend(view2), KE::begin(view2), args...);
ASSERT_EQ(r, KE::end(view2));
VerifyData()(view1, view2, args...);
}
{
fill_zero(view_dest);
auto r = KE::inclusive_scan(exespace(), view_from, view_dest, bop, args...);
ASSERT_EQ(r, KE::end(view_dest));
verify_data(view_from, view_dest, bop, args...);
fill_view(view2, name);
auto r = KE::inclusive_scan(exespace(), view2, view2, args...);
ASSERT_EQ(r, KE::end(view2));
VerifyData()(view1, view2, args...);
}
{
fill_zero(view_dest);
auto r = KE::inclusive_scan("label", exespace(), view_from, view_dest, bop,
args...);
ASSERT_EQ(r, KE::end(view_dest));
verify_data(view_from, view_dest, bop, args...);
fill_view(view2, name);
auto r = KE::inclusive_scan("label", exespace(), view2, view2, args...);
ASSERT_EQ(r, KE::end(view2));
VerifyData()(view1, view2, args...);
}
Kokkos::fence();
@ -318,27 +318,35 @@ void run_inclusive_scan_all_scenarios() {
{"medium-a", 313}, {"medium-b", 1103}, {"large", 10513}};
for (const auto& it : scenarios) {
run_single_scenario_default_op<Tag, ValueType>(it);
run_single_scenario<Tag, ValueType>(it);
run_single_scenario_inplace<Tag, ValueType>(it);
#if !defined KOKKOS_ENABLE_OPENMPTARGET
// the sum custom op is always run
using sum_binary_op = SumFunctor<ValueType>;
sum_binary_op sbop;
run_single_scenario_custom_op<Tag, ValueType>(it, sbop);
run_single_scenario_custom_op<Tag, ValueType>(it, sbop, ValueType{0});
run_single_scenario_custom_op<Tag, ValueType>(it, sbop, ValueType{1});
run_single_scenario_custom_op<Tag, ValueType>(it, sbop, ValueType{-2});
run_single_scenario_custom_op<Tag, ValueType>(it, sbop, ValueType{3});
run_single_scenario<Tag, ValueType>(it, sbop);
run_single_scenario<Tag, ValueType>(it, sbop, ValueType{0});
run_single_scenario<Tag, ValueType>(it, sbop, ValueType{1});
run_single_scenario<Tag, ValueType>(it, sbop, ValueType{-2});
run_single_scenario<Tag, ValueType>(it, sbop, ValueType{3});
run_single_scenario_inplace<Tag, ValueType>(it, sbop, ValueType{0});
run_single_scenario_inplace<Tag, ValueType>(it, sbop, ValueType{-2});
// custom multiply only for small views to avoid overflows
if (it.first == "small-a" || it.first == "small-b") {
using mult_binary_op = MultiplyFunctor<ValueType>;
mult_binary_op mbop;
run_single_scenario_custom_op<Tag, ValueType>(it, mbop);
run_single_scenario_custom_op<Tag, ValueType>(it, mbop, ValueType{0});
run_single_scenario_custom_op<Tag, ValueType>(it, mbop, ValueType{1});
run_single_scenario_custom_op<Tag, ValueType>(it, mbop, ValueType{-2});
run_single_scenario_custom_op<Tag, ValueType>(it, mbop, ValueType{3});
run_single_scenario<Tag, ValueType>(it, mbop);
run_single_scenario<Tag, ValueType>(it, mbop, ValueType{0});
run_single_scenario<Tag, ValueType>(it, mbop, ValueType{1});
run_single_scenario<Tag, ValueType>(it, mbop, ValueType{-2});
run_single_scenario<Tag, ValueType>(it, mbop, ValueType{3});
run_single_scenario_inplace<Tag, ValueType>(it, mbop);
run_single_scenario_inplace<Tag, ValueType>(it, mbop, ValueType{0});
run_single_scenario_inplace<Tag, ValueType>(it, mbop, ValueType{-2});
}
#endif
}

7
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsIsSorted.cpp
View File

@ -146,7 +146,7 @@ void run_single_scenario(const InfoType& scenario_info) {
resultsA[3] = KE::is_sorted("label", exespace(), view);
const auto allA = std::all_of(resultsA.cbegin(), resultsA.cend(),
[=](bool v) { return v == gold; });
EXPECT_TRUE(allA);
EXPECT_TRUE(allA) << name << ", " << view_tag_to_string(Tag{});
#if !defined KOKKOS_ENABLE_OPENMPTARGET
CustomLessThanComparator<ValueType, ValueType> comp;
@ -159,7 +159,7 @@ void run_single_scenario(const InfoType& scenario_info) {
resultsB[3] = KE::is_sorted("label", exespace(), view, comp);
const auto allB = std::all_of(resultsB.cbegin(), resultsB.cend(),
[=](bool v) { return v == gold; });
EXPECT_TRUE(allB);
EXPECT_TRUE(allB) << name << ", " << view_tag_to_string(Tag{});
#endif
Kokkos::fence();
@ -173,9 +173,6 @@ void run_is_sorted_all_scenarios() {
{"medium-a", 1003}, {"medium-b", 1003}, {"large-a", 101513},
{"large-b", 101513}};
std::cout << "is_sorted: " << view_tag_to_string(Tag{})
<< ", all overloads \n";
for (const auto& it : scenarios) {
run_single_scenario<Tag, ValueType>(it);
}

19
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);
ASSERT_EQ(r1, gold);
ASSERT_EQ(r2, gold);
ASSERT_EQ(r3, gold);
ASSERT_EQ(r4, gold);
ASSERT_EQ(r1, gold) << name << ", " << view_tag_to_string(Tag{});
ASSERT_EQ(r2, gold) << name << ", " << view_tag_to_string(Tag{});
ASSERT_EQ(r3, gold) << name << ", " << view_tag_to_string(Tag{});
ASSERT_EQ(r4, gold) << name << ", " << view_tag_to_string(Tag{});
#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
ASSERT_EQ(r1, gold);
ASSERT_EQ(r2, gold);
ASSERT_EQ(r3, gold);
ASSERT_EQ(r4, gold);
ASSERT_EQ(r1, gold) << name << ", " << view_tag_to_string(Tag{});
ASSERT_EQ(r2, gold) << name << ", " << view_tag_to_string(Tag{});
ASSERT_EQ(r3, gold) << name << ", " << view_tag_to_string(Tag{});
ASSERT_EQ(r4, gold) << name << ", " << view_tag_to_string(Tag{});
Kokkos::fence();
}
@ -176,9 +176,6 @@ void run_is_sorted_until_all_scenarios() {
{"medium-a", 1003}, {"medium-b", 1003}, {"large-a", 101513},
{"large-b", 101513}};
std::cout << "is_sorted_until: " << view_tag_to_string(Tag{})
<< ", all overloads \n";
for (const auto& it : scenarios) {
run_single_scenario<Tag, ValueType>(it);
}

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

@ -48,7 +48,7 @@ struct MyMovableType {
TEST(std_algorithms_mod_ops_test, move) {
MyMovableType a;
using move_t = decltype(std::move(a));
static_assert(std::is_rvalue_reference<move_t>::value, "");
static_assert(std::is_rvalue_reference<move_t>::value);
// move constr
MyMovableType b(std::move(a));
@ -70,7 +70,7 @@ struct StdAlgoModSeqOpsTestMove {
void operator()(const int index) const {
typename ViewType::value_type a{11};
using move_t = decltype(std::move(a));
static_assert(std::is_rvalue_reference<move_t>::value, "");
static_assert(std::is_rvalue_reference<move_t>::value);
m_view(index) = std::move(a);
}
@ -89,50 +89,6 @@ TEST(std_algorithms_mod_ops_test, move_within_parfor) {
}
}
// ------------
// swap
// ------------
TEST(std_algorithms_mod_ops_test, swap) {
{
int a = 1;
int b = 2;
KE::swap(a, b);
ASSERT_EQ(a, 2);
ASSERT_EQ(b, 1);
}
{
double a = 3.;
double b = 1.;
KE::swap(a, b);
EXPECT_DOUBLE_EQ(a, 1.);
EXPECT_DOUBLE_EQ(b, 3.);
}
}
template <class ViewType>
struct StdAlgoModSeqOpsTestSwap {
ViewType m_view;
KOKKOS_INLINE_FUNCTION
void operator()(const int index) const {
typename ViewType::value_type newval{11};
KE::swap(m_view(index), newval);
}
StdAlgoModSeqOpsTestSwap(ViewType aIn) : m_view(aIn) {}
};
TEST(std_algorithms_mod_ops_test, swap_within_parfor) {
auto a = create_view<double>(stdalgos::DynamicTag{}, 10, "a");
StdAlgoModSeqOpsTestSwap<decltype(a)> fnc(a);
Kokkos::parallel_for(a.extent(0), fnc);
auto a_h = Kokkos::create_mirror_view_and_copy(Kokkos::HostSpace(), a);
for (std::size_t i = 0; i < a.extent(0); ++i) {
EXPECT_DOUBLE_EQ(a_h(0), 11.);
}
}
// ------------
// iter_swap
// ------------

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

@ -110,11 +110,9 @@ void verify_data(const std::string& name, ResultType my_result,
ViewTypeDestFalse view_dest_false, PredType pred) {
using value_type = typename ViewTypeFrom::value_type;
static_assert(
std::is_same<value_type, typename ViewTypeDestTrue::value_type>::value,
"");
std::is_same<value_type, typename ViewTypeDestTrue::value_type>::value);
static_assert(
std::is_same<value_type, typename ViewTypeDestFalse::value_type>::value,
"");
std::is_same<value_type, typename ViewTypeDestFalse::value_type>::value);
const std::size_t ext = view_from.extent(0);

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

@ -166,6 +166,10 @@ void run_all_scenarios() {
}
TEST(std_algorithms_copy_if_team_test, test) {
// FIXME_OPENMPTARGET
#if defined(KOKKOS_ENABLE_OPENMPTARGET) && defined(KOKKOS_ARCH_INTEL_GPU)
GTEST_SKIP() << "the test is known to fail with OpenMPTarget on Intel GPUs";
#endif
run_all_scenarios<DynamicTag, double>();
run_all_scenarios<StridedTwoRowsTag, int>();
run_all_scenarios<StridedThreeRowsTag, unsigned>();

43
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsTeamExclusiveScan.cpp
View File

@ -121,7 +121,9 @@ struct TestFunctorA {
}
};
template <class LayoutTag, class ValueType>
struct InPlace {};
template <class LayoutTag, class ValueType, class InPlaceOrVoid = void>
void test_A(std::size_t numTeams, std::size_t numCols, int apiId) {
/* description:
use a rank-2 view randomly filled with values,
@ -147,9 +149,6 @@ void test_A(std::size_t numTeams, std::size_t numCols, int apiId) {
using space_t = Kokkos::DefaultExecutionSpace;
Kokkos::TeamPolicy<space_t> policy(numTeams, Kokkos::AUTO());
// create the destination view
Kokkos::View<ValueType**> destView("destView", numTeams, numCols);
// exclusive_scan returns an iterator so to verify that it is correct
// each team stores the distance of the returned iterator from the beginning
// of the interval that team operates on and then we check that these
@ -168,12 +167,19 @@ void test_A(std::size_t numTeams, std::size_t numCols, int apiId) {
rand_pool pool(lowerBound * upperBound);
Kokkos::fill_random(initValuesView_h, pool, lowerBound, upperBound);
// use CTAD for functor
auto initValuesView =
Kokkos::create_mirror_view_and_copy(space_t(), initValuesView_h);
TestFunctorA fnc(sourceView, destView, distancesView, intraTeamSentinelView,
initValuesView, binaryOp, apiId);
Kokkos::parallel_for(policy, fnc);
Kokkos::View<ValueType**> destView("destView", numTeams, numCols);
if constexpr (std::is_same_v<InPlaceOrVoid, InPlace>) {
TestFunctorA fnc(sourceView, sourceView, distancesView,
intraTeamSentinelView, initValuesView, binaryOp, apiId);
Kokkos::parallel_for(policy, fnc);
} else {
TestFunctorA fnc(sourceView, destView, distancesView, intraTeamSentinelView,
initValuesView, binaryOp, apiId);
Kokkos::parallel_for(policy, fnc);
}
// -----------------------------------------------
// run cpp-std kernel and check
@ -223,11 +229,16 @@ void test_A(std::size_t numTeams, std::size_t numCols, int apiId) {
#undef exclusive_scan
}
auto dataViewAfterOp_h = create_host_space_copy(destView);
expect_equal_host_views(stdDestView, dataViewAfterOp_h);
if constexpr (std::is_same_v<InPlaceOrVoid, InPlace>) {
auto dataViewAfterOp_h = create_host_space_copy(sourceView);
expect_equal_host_views(stdDestView, dataViewAfterOp_h);
} else {
auto dataViewAfterOp_h = create_host_space_copy(destView);
expect_equal_host_views(stdDestView, dataViewAfterOp_h);
}
}
template <class LayoutTag, class ValueType>
template <class LayoutTag, class ValueType, class InPlaceOrVoid = void>
void run_all_scenarios() {
for (int numTeams : teamSizesToTest) {
for (const auto& numCols : {0, 1, 2, 13, 101, 1444, 8153}) {
@ -236,16 +247,24 @@ void run_all_scenarios() {
#else
for (int apiId : {0, 1}) {
#endif
test_A<LayoutTag, ValueType>(numTeams, numCols, apiId);
test_A<LayoutTag, ValueType, InPlaceOrVoid>(numTeams, numCols, apiId);
}
}
}
}
TEST(std_algorithms_exclusive_scan_team_test, test) {
// FIXME_OPENMPTARGET
#if defined(KOKKOS_ENABLE_OPENMPTARGET) && defined(KOKKOS_ARCH_INTEL_GPU)
GTEST_SKIP() << "the test is known to fail with OpenMPTarget on Intel GPUs";
#endif
run_all_scenarios<DynamicTag, double>();
run_all_scenarios<StridedTwoRowsTag, int>();
run_all_scenarios<StridedThreeRowsTag, unsigned>();
run_all_scenarios<DynamicTag, double, InPlace>();
run_all_scenarios<StridedTwoRowsTag, int, InPlace>();
run_all_scenarios<StridedThreeRowsTag, unsigned, InPlace>();
}
} // namespace TeamExclusiveScan

44
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsTeamInclusiveScan.cpp
View File

@ -139,7 +139,9 @@ struct TestFunctorA {
}
};
template <class LayoutTag, class ValueType>
struct InPlace {};
template <class LayoutTag, class ValueType, class InPlaceOrVoid = void>
void test_A(std::size_t numTeams, std::size_t numCols, int apiId) {
/* description:
use a rank-2 view randomly filled with values,
@ -165,9 +167,6 @@ void test_A(std::size_t numTeams, std::size_t numCols, int apiId) {
using space_t = Kokkos::DefaultExecutionSpace;
Kokkos::TeamPolicy<space_t> policy(numTeams, Kokkos::AUTO());
// create the destination view
Kokkos::View<ValueType**> destView("destView", numTeams, numCols);
// inclusive_scan returns an iterator so to verify that it is correct
// each team stores the distance of the returned iterator from the beginning
// of the interval that team operates on and then we check that these
@ -186,12 +185,20 @@ void test_A(std::size_t numTeams, std::size_t numCols, int apiId) {
rand_pool pool(lowerBound * upperBound);
Kokkos::fill_random(initValuesView_h, pool, lowerBound, upperBound);
// use CTAD for functor
auto initValuesView =
Kokkos::create_mirror_view_and_copy(space_t(), initValuesView_h);
TestFunctorA fnc(sourceView, destView, distancesView, intraTeamSentinelView,
initValuesView, binaryOp, apiId);
Kokkos::parallel_for(policy, fnc);
// create the destination view
Kokkos::View<ValueType**> destView("destView", numTeams, numCols);
if constexpr (std::is_same_v<InPlaceOrVoid, InPlace>) {
TestFunctorA fnc(sourceView, sourceView, distancesView,
intraTeamSentinelView, initValuesView, binaryOp, apiId);
Kokkos::parallel_for(policy, fnc);
} else {
TestFunctorA fnc(sourceView, destView, distancesView, intraTeamSentinelView,
initValuesView, binaryOp, apiId);
Kokkos::parallel_for(policy, fnc);
}
// -----------------------------------------------
// run cpp-std kernel and check
@ -251,25 +258,38 @@ void test_A(std::size_t numTeams, std::size_t numCols, int apiId) {
#undef inclusive_scan
}
auto dataViewAfterOp_h = create_host_space_copy(destView);
expect_equal_host_views(stdDestView, dataViewAfterOp_h);
if constexpr (std::is_same_v<InPlaceOrVoid, InPlace>) {
auto dataViewAfterOp_h = create_host_space_copy(sourceView);
expect_equal_host_views(stdDestView, dataViewAfterOp_h);
} else {
auto dataViewAfterOp_h = create_host_space_copy(destView);
expect_equal_host_views(stdDestView, dataViewAfterOp_h);
}
}
template <class LayoutTag, class ValueType>
template <class LayoutTag, class ValueType, class InPlaceOrVoid = void>
void run_all_scenarios() {
for (int numTeams : teamSizesToTest) {
for (const auto& numCols : {0, 1, 2, 13, 101, 1444, 8153}) {
for (int apiId : {0, 1, 2, 3, 4, 5}) {
test_A<LayoutTag, ValueType>(numTeams, numCols, apiId);
test_A<LayoutTag, ValueType, InPlaceOrVoid>(numTeams, numCols, apiId);
}
}
}
}
TEST(std_algorithms_inclusive_scan_team_test, test) {
// FIXME_OPENMPTARGET
#if defined(KOKKOS_ENABLE_OPENMPTARGET) && defined(KOKKOS_ARCH_INTEL_GPU)
GTEST_SKIP() << "the test is known to fail with OpenMPTarget on Intel GPUs";
#endif
run_all_scenarios<DynamicTag, double>();
run_all_scenarios<StridedTwoRowsTag, int>();
run_all_scenarios<StridedThreeRowsTag, unsigned>();
run_all_scenarios<DynamicTag, double, InPlace>();
run_all_scenarios<StridedTwoRowsTag, int, InPlace>();
run_all_scenarios<StridedThreeRowsTag, unsigned, InPlace>();
}
} // namespace TeamInclusiveScan

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

@ -212,6 +212,10 @@ void run_all_scenarios() {
}
TEST(std_algorithms_remove_copy_team_test, test) {
// FIXME_OPENMPTARGET
#if defined(KOKKOS_ENABLE_OPENMPTARGET) && defined(KOKKOS_ARCH_INTEL_GPU)
GTEST_SKIP() << "the test is known to fail with OpenMPTarget on Intel GPUs";
#endif
run_all_scenarios<DynamicTag, double>();
run_all_scenarios<StridedTwoRowsTag, int>();
run_all_scenarios<StridedThreeRowsTag, unsigned>();

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

@ -168,6 +168,10 @@ void run_all_scenarios() {
}
TEST(std_algorithms_remove_copy_if_team_test, test) {
// FIXME_OPENMPTARGET
#if defined(KOKKOS_ENABLE_OPENMPTARGET) && defined(KOKKOS_ARCH_INTEL_GPU)
GTEST_SKIP() << "the test is known to fail with OpenMPTarget on Intel GPUs";
#endif
run_all_scenarios<DynamicTag, double>();
run_all_scenarios<StridedTwoRowsTag, int>();
run_all_scenarios<StridedThreeRowsTag, unsigned>();

41
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsTeamTransformExclusiveScan.cpp
View File

@ -108,7 +108,9 @@ struct TestFunctorA {
}
};
template <class LayoutTag, class ValueType>
struct InPlace {};
template <class LayoutTag, class ValueType, class InPlaceOrVoid = void>
void test_A(std::size_t numTeams, std::size_t numCols, int apiId) {
/* description:
use a rank-2 view randomly filled with values,
@ -134,9 +136,6 @@ void test_A(std::size_t numTeams, std::size_t numCols, int apiId) {
using space_t = Kokkos::DefaultExecutionSpace;
Kokkos::TeamPolicy<space_t> policy(numTeams, Kokkos::AUTO());
// create the destination view
Kokkos::View<ValueType**> destView("destView", numTeams, numCols);
// tranform_exclusive_scan returns an iterator so to verify that it is correct
// each team stores the distance of the returned iterator from the beginning
// of the interval that team operates on and then we check that these
@ -156,12 +155,21 @@ void test_A(std::size_t numTeams, std::size_t numCols, int apiId) {
rand_pool pool(lowerBound * upperBound);
Kokkos::fill_random(initValuesView_h, pool, lowerBound, upperBound);
// use CTAD for functor
auto initValuesView =
Kokkos::create_mirror_view_and_copy(space_t(), initValuesView_h);
TestFunctorA fnc(sourceView, destView, distancesView, intraTeamSentinelView,
initValuesView, binaryOp, unaryOp, apiId);
Kokkos::parallel_for(policy, fnc);
// create the destination view
Kokkos::View<ValueType**> destView("destView", numTeams, numCols);
if constexpr (std::is_same_v<InPlaceOrVoid, InPlace>) {
TestFunctorA fnc(sourceView, sourceView, distancesView,
intraTeamSentinelView, initValuesView, binaryOp, unaryOp,
apiId);
Kokkos::parallel_for(policy, fnc);
} else {
TestFunctorA fnc(sourceView, destView, distancesView, intraTeamSentinelView,
initValuesView, binaryOp, unaryOp, apiId);
Kokkos::parallel_for(policy, fnc);
}
// -----------------------------------------------
// run cpp-std kernel and check
@ -200,16 +208,21 @@ void test_A(std::size_t numTeams, std::size_t numCols, int apiId) {
#undef transform_exclusive_scan
}
auto dataViewAfterOp_h = create_host_space_copy(destView);
expect_equal_host_views(stdDestView, dataViewAfterOp_h);
if constexpr (std::is_same_v<InPlaceOrVoid, InPlace>) {
auto dataViewAfterOp_h = create_host_space_copy(sourceView);
expect_equal_host_views(stdDestView, dataViewAfterOp_h);
} else {
auto dataViewAfterOp_h = create_host_space_copy(destView);
expect_equal_host_views(stdDestView, dataViewAfterOp_h);
}
}
template <class LayoutTag, class ValueType>
template <class LayoutTag, class ValueType, class InPlaceOrVoid = void>
void run_all_scenarios() {
for (int numTeams : teamSizesToTest) {
for (const auto& numCols : {0, 1, 2, 13, 101, 1444, 8153}) {
for (int apiId : {0, 1}) {
test_A<LayoutTag, ValueType>(numTeams, numCols, apiId);
test_A<LayoutTag, ValueType, InPlaceOrVoid>(numTeams, numCols, apiId);
}
}
}
@ -219,6 +232,10 @@ TEST(std_algorithms_transform_exclusive_scan_team_test, test) {
run_all_scenarios<DynamicTag, double>();
run_all_scenarios<StridedTwoRowsTag, int>();
run_all_scenarios<StridedThreeRowsTag, unsigned>();
run_all_scenarios<DynamicTag, double, InPlace>();
run_all_scenarios<StridedTwoRowsTag, int, InPlace>();
run_all_scenarios<StridedThreeRowsTag, unsigned, InPlace>();
}
} // namespace TeamTransformExclusiveScan

41
lib/kokkos/algorithms/unit_tests/TestStdAlgorithmsTeamTransformInclusiveScan.cpp
View File

@ -131,7 +131,9 @@ struct TestFunctorA {
}
};
template <class LayoutTag, class ValueType>
struct InPlace {};
template <class LayoutTag, class ValueType, class InPlaceOrVoid = void>
void test_A(std::size_t numTeams, std::size_t numCols, int apiId) {
/* description:
use a rank-2 view randomly filled with values,
@ -157,9 +159,6 @@ void test_A(std::size_t numTeams, std::size_t numCols, int apiId) {
using space_t = Kokkos::DefaultExecutionSpace;
Kokkos::TeamPolicy<space_t> policy(numTeams, Kokkos::AUTO());
// create the destination view
Kokkos::View<ValueType**> destView("destView", numTeams, numCols);
// tranform_inclusive_scan returns an iterator so to verify that it is correct
// each team stores the distance of the returned iterator from the beginning
// of the interval that team operates on and then we check that these
@ -179,12 +178,21 @@ void test_A(std::size_t numTeams, std::size_t numCols, int apiId) {
rand_pool pool(lowerBound * upperBound);
Kokkos::fill_random(initValuesView_h, pool, lowerBound, upperBound);
// use CTAD for functor
auto initValuesView =
Kokkos::create_mirror_view_and_copy(space_t(), initValuesView_h);
TestFunctorA fnc(sourceView, destView, distancesView, intraTeamSentinelView,
initValuesView, binaryOp, unaryOp, apiId);
Kokkos::parallel_for(policy, fnc);
// create the destination view
Kokkos::View<ValueType**> destView("destView", numTeams, numCols);
if constexpr (std::is_same_v<InPlaceOrVoid, InPlace>) {
TestFunctorA fnc(sourceView, sourceView, distancesView,
intraTeamSentinelView, initValuesView, binaryOp, unaryOp,
apiId);
Kokkos::parallel_for(policy, fnc);
} else {
TestFunctorA fnc(sourceView, destView, distancesView, intraTeamSentinelView,
initValuesView, binaryOp, unaryOp, apiId);
Kokkos::parallel_for(policy, fnc);
}
// -----------------------------------------------
// run cpp-std kernel and check
@ -236,16 +244,21 @@ void test_A(std::size_t numTeams, std::size_t numCols, int apiId) {
}
#undef transform_inclusive_scan
auto dataViewAfterOp_h = create_host_space_copy(destView);
expect_equal_host_views(stdDestView, dataViewAfterOp_h);
if constexpr (std::is_same_v<InPlaceOrVoid, InPlace>) {
auto dataViewAfterOp_h = create_host_space_copy(sourceView);
expect_equal_host_views(stdDestView, dataViewAfterOp_h);
} else {
auto dataViewAfterOp_h = create_host_space_copy(destView);
expect_equal_host_views(stdDestView, dataViewAfterOp_h);
}
}
template <class LayoutTag, class ValueType>
template <class LayoutTag, class ValueType, class InPlaceOrVoid = void>
void run_all_scenarios() {
for (int numTeams : teamSizesToTest) {
for (const auto& numCols : {0, 1, 2, 13, 101, 1444, 8153}) {
for (int apiId : {0, 1, 2, 3}) {
test_A<LayoutTag, ValueType>(numTeams, numCols, apiId);
test_A<LayoutTag, ValueType, InPlaceOrVoid>(numTeams, numCols, apiId);
}
}
}
@ -255,6 +268,10 @@ TEST(std_algorithms_transform_inclusive_scan_team_test, test) {
run_all_scenarios<DynamicTag, double>();
run_all_scenarios<StridedTwoRowsTag, int>();
run_all_scenarios<StridedThreeRowsTag, unsigned>();
run_all_scenarios<DynamicTag, double, InPlace>();
run_all_scenarios<StridedTwoRowsTag, int, InPlace>();
run_all_scenarios<StridedThreeRowsTag, unsigned, InPlace>();
}
} // namespace TeamTransformInclusiveScan

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

@ -186,6 +186,10 @@ void run_all_scenarios() {
}
TEST(std_algorithms_unique_copy_team_test, test) {
// FIXME_OPENMPTARGET
#if defined(KOKKOS_ENABLE_OPENMPTARGET) && defined(KOKKOS_ARCH_INTEL_GPU)
GTEST_SKIP() << "the test is known to fail with OpenMPTarget on Intel GPUs";
#endif
run_all_scenarios<DynamicTag, int>();
run_all_scenarios<StridedTwoRowsTag, int>();
run_all_scenarios<StridedThreeRowsTag, int>();

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

@ -16,6 +16,7 @@
#include <TestStdAlgorithmsCommon.hpp>
#include <utility>
#include <iomanip>
namespace Test {
namespace stdalgos {
@ -160,24 +161,15 @@ void verify_data(ViewType1 data_view, // contains data
create_mirror_view_and_copy(Kokkos::HostSpace(), test_view_dc);
if (test_view_h.extent(0) > 0) {
for (std::size_t i = 0; i < test_view_h.extent(0); ++i) {
// std::cout << i << " " << std::setprecision(15) << data_view_h(i) << " "
// << 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) {
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) {
std::cout << i << " " << std::setprecision(15) << data_view_h(i)
<< " " << gold_h(i) << " " << test_view_h(i) << " "
<< std::abs(gold_h(i) - test_view_h(i)) << std::endl;
}
EXPECT_LT(error, 1e-10);
ASSERT_LT(error, 1e-10) << i << " " << std::setprecision(15) << error
<< static_cast<double>(test_view_h(i)) << " "
<< static_cast<double>(gold_h(i));
}
}
// std::cout << " last el: " << test_view_h(test_view_h.extent(0)-1) <<
// std::endl;
}
}
@ -205,17 +197,13 @@ void run_single_scenario(const InfoType& scenario_info, ValueType init_value,
BinaryOp bop, UnaryOp uop) {
const auto name = std::get<0>(scenario_info);
const std::size_t view_ext = std::get<1>(scenario_info);
// std::cout << "transform_exclusive_scan custom op: " << name << ", "
// << view_tag_to_string(Tag{}) << ", "
// << value_type_to_string(ValueType()) << ", "
// << "init = " << init_value << std::endl;
auto view_dest =
create_view<ValueType>(Tag{}, view_ext, "transform_exclusive_scan");
auto view_from =
create_view<ValueType>(Tag{}, view_ext, "transform_exclusive_scan");
auto view_from = create_view<ValueType>(Tag{}, view_ext,
"transform_exclusive_scan_view_from");
fill_view(view_from, name);
auto view_dest = create_view<ValueType>(Tag{}, view_ext,
"transform_exclusive_scan_view_dest");
{
fill_zero(view_dest);
auto r = KE::transform_exclusive_scan(
@ -253,6 +241,65 @@ void run_single_scenario(const InfoType& scenario_info, ValueType init_value,
Kokkos::fence();
}
template <class Tag, class ValueType, class InfoType, class BinaryOp,
class UnaryOp>
void run_single_scenario_inplace(const InfoType& scenario_info,
ValueType init_value, BinaryOp bop,
UnaryOp uop) {
const auto name = std::get<0>(scenario_info);
const std::size_t view_ext = std::get<1>(scenario_info);
// since here we call the in-place operation, we need to use two views:
// view1: filled according to what the scenario asks for and is not modified
// view2: filled according to what the scenario asks for and used for the
// in-place op Therefore, after the op is done, view2 should contain the
// result of doing exclusive scan NOTE: view2 is filled below every time
// because the algorithm acts in place
auto view1 =
create_view<ValueType>(Tag{}, view_ext, "transform_exclusive_scan_view1");
fill_view(view1, name);
auto view2 =
create_view<ValueType>(Tag{}, view_ext, "transform_exclusive_scan_view2");
{
fill_view(view2, name);
auto r = KE::transform_exclusive_scan(exespace(), KE::cbegin(view2),
KE::cend(view2), KE::begin(view2),
init_value, bop, uop);
ASSERT_EQ(r, KE::end(view2));
verify_data(view1, view2, init_value, bop, uop);
}
{
fill_view(view2, name);
auto r = KE::transform_exclusive_scan(
"label", exespace(), KE::cbegin(view2), KE::cend(view2),
KE::begin(view2), init_value, bop, uop);
ASSERT_EQ(r, KE::end(view2));
verify_data(view1, view2, init_value, bop, uop);
}
{
fill_view(view2, name);
auto r = KE::transform_exclusive_scan(exespace(), view2, view2, init_value,
bop, uop);
ASSERT_EQ(r, KE::end(view2));
verify_data(view1, view2, init_value, bop, uop);
}
{
fill_view(view2, name);
auto r = KE::transform_exclusive_scan("label", exespace(), view2, view2,
init_value, bop, uop);
ASSERT_EQ(r, KE::end(view2));
verify_data(view1, view2, init_value, bop, uop);
}
Kokkos::fence();
}
template <class Tag, class ValueType>
void run_all_scenarios() {
const std::map<std::string, std::size_t> scenarios = {
@ -267,6 +314,11 @@ void run_all_scenarios() {
run_single_scenario<Tag, ValueType>(it, ValueType{1}, bop_t(), uop_t());
run_single_scenario<Tag, ValueType>(it, ValueType{-2}, bop_t(), uop_t());
run_single_scenario<Tag, ValueType>(it, ValueType{3}, bop_t(), uop_t());
run_single_scenario_inplace<Tag, ValueType>(it, ValueType{0}, bop_t(),
uop_t());
run_single_scenario_inplace<Tag, ValueType>(it, ValueType{-2}, bop_t(),
uop_t());
}
}

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

@ -16,6 +16,7 @@
#include <TestStdAlgorithmsCommon.hpp>
#include <utility>
#include <iomanip>
namespace Test {
namespace stdalgos {
@ -172,24 +173,15 @@ void verify_data(ViewType1 data_view, // contains data
create_mirror_view_and_copy(Kokkos::HostSpace(), test_view_dc);
if (test_view_h.extent(0) > 0) {
for (std::size_t i = 0; i < test_view_h.extent(0); ++i) {
// std::cout << i << " " << std::setprecision(15) << data_view_h(i) << " "
// << 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) {
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) {
std::cout << i << " " << std::setprecision(15) << data_view_h(i)
<< " " << gold_h(i) << " " << test_view_h(i) << " "
<< std::abs(gold_h(i) - test_view_h(i)) << std::endl;
}
EXPECT_LT(error, 1e-10);
ASSERT_LT(error, 1e-10) << i << " " << std::setprecision(15) << error
<< static_cast<double>(test_view_h(i)) << " "
<< static_cast<double>(gold_h(i));
}
}
// std::cout << " last el: " << test_view_h(test_view_h.extent(0)-1) <<
// std::endl;
}
}
@ -210,30 +202,11 @@ struct SumBinaryFunctor {
std::string value_type_to_string(int) { return "int"; }
std::string value_type_to_string(double) { return "double"; }
template <class Tag, class BopT, class UopT>
void print_scenario_details(const std::string& name, BopT bop, UopT uop) {
(void)bop;
(void)uop;
std::cout << "transform_inclusive_scan: " << name << ", "
<< view_tag_to_string(Tag{}) << std::endl;
}
template <class Tag, class BopT, class UopT, class ValueType>
void print_scenario_details(const std::string& name, BopT bop, UopT uop,
ValueType init_value) {
(void)bop;
(void)uop;
std::cout << "transform_inclusive_scan: " << name << ", "
<< view_tag_to_string(Tag{}) << ", "
<< "init = " << init_value << std::endl;
}
template <class Tag, class ValueType, class InfoType, class... Args>
void run_single_scenario(const InfoType& scenario_info,
Args... args /* by value on purpose*/) {
const auto name = std::get<0>(scenario_info);
const std::size_t view_ext = std::get<1>(scenario_info);
// print_scenario_details<Tag>(name, args...);
auto view_dest =
create_view<ValueType>(Tag{}, view_ext, "transform_inclusive_scan");
@ -278,6 +251,63 @@ void run_single_scenario(const InfoType& scenario_info,
Kokkos::fence();
}
template <class Tag, class ValueType, class InfoType, class... Args>
void run_single_scenario_inplace(const InfoType& scenario_info,
Args... args /* by value on purpose*/) {
const auto name = std::get<0>(scenario_info);
const std::size_t view_ext = std::get<1>(scenario_info);
// since here we call the in-place operation, we need to use two views:
// view1: filled according to scenario and is not modified
// view2: filled according scenario and used for the in-place op
// Therefore, after the op is done, view_2 should contain the
// result of doing exclusive scan.
// NOTE: view2 must be filled before every call to the algorithm
// because the algorithm acts in place
auto view_1 = create_view<ValueType>(Tag{}, view_ext,
"transform_inclusive_scan_view_1");
fill_view(view_1, name);
auto view_2 = create_view<ValueType>(Tag{}, view_ext,
"transform_inclusive_scan_view_2");
{
fill_view(view_2, name);
auto r = KE::transform_inclusive_scan(exespace(), KE::cbegin(view_2),
KE::cend(view_2), KE::begin(view_2),
args...);
ASSERT_EQ(r, KE::end(view_2));
verify_data(view_1, view_2, args...);
}
{
fill_view(view_2, name);
auto r = KE::transform_inclusive_scan("label", exespace(),
KE::cbegin(view_2), KE::cend(view_2),
KE::begin(view_2), args...);
ASSERT_EQ(r, KE::end(view_2));
verify_data(view_1, view_2, args...);
}
{
fill_view(view_2, name);
auto r = KE::transform_inclusive_scan(exespace(), view_2, view_2, args...);
ASSERT_EQ(r, KE::end(view_2));
verify_data(view_1, view_2, args...);
}
{
fill_view(view_2, name);
auto r = KE::transform_inclusive_scan("label", exespace(), view_2, view_2,
args...);
ASSERT_EQ(r, KE::end(view_2));
verify_data(view_1, view_2, args...);
}
Kokkos::fence();
}
template <class Tag, class ValueType>
void run_all_scenarios() {
const std::map<std::string, std::size_t> scenarios = {
@ -294,15 +324,23 @@ void run_all_scenarios() {
run_single_scenario<Tag, ValueType>(it, bop_t(), uop_t(), ValueType{2});
run_single_scenario<Tag, ValueType>(it, bop_t(), uop_t(), ValueType{-1});
run_single_scenario<Tag, ValueType>(it, bop_t(), uop_t(), ValueType{-2});
run_single_scenario_inplace<Tag, ValueType>(it, bop_t(), uop_t());
run_single_scenario_inplace<Tag, ValueType>(it, bop_t(), uop_t(),
ValueType{0});
run_single_scenario_inplace<Tag, ValueType>(it, bop_t(), uop_t(),
ValueType{2});
run_single_scenario_inplace<Tag, ValueType>(it, bop_t(), uop_t(),
ValueType{-2});
}
}
#if !defined KOKKOS_ENABLE_OPENMPTARGET
TEST(std_algorithms_numeric_ops_test, transform_inclusive_scan) {
run_all_scenarios<DynamicTag, double>();
// run_all_scenarios<StridedThreeTag, double>();
// run_all_scenarios<DynamicTag, int>();
// run_all_scenarios<StridedThreeTag, int>();
run_all_scenarios<StridedThreeTag, double>();
run_all_scenarios<DynamicTag, int>();
run_all_scenarios<StridedThreeTag, int>();
}
#endif

32
lib/kokkos/algorithms/unit_tests/TestStdReducers.cpp
View File

@ -83,9 +83,6 @@ auto run_min_or_max_test(ViewType view, StdReducersTestEnumOrder enValue) {
static_assert(std::is_same<ExeSpace, Kokkos::HostSpace>::value,
"test is only enabled for HostSpace");
std::cout << "checking reduction with order: " << order_to_string(enValue)
<< "\n";
using view_value_type = typename ViewType::value_type;
using reducer_type = std::conditional_t<
(flag == 0), Kokkos::MaxFirstLoc<view_value_type, IndexType, ExeSpace>,
@ -132,18 +129,24 @@ TEST(std_algorithms_reducers, max_first_loc) {
const auto pair1 = run_min_or_max_test<0, hostspace, index_type>(
view_h, StdReducersTestEnumOrder::LeftToRight);
ASSERT_EQ(pair1.first, gold_value);
ASSERT_EQ(pair1.second, gold_location);
ASSERT_EQ(pair1.first, gold_value)
<< order_to_string(StdReducersTestEnumOrder::LeftToRight);
ASSERT_EQ(pair1.second, gold_location)
<< order_to_string(StdReducersTestEnumOrder::LeftToRight);
const auto pair2 = run_min_or_max_test<0, hostspace, index_type>(
view_h, StdReducersTestEnumOrder::RightToLeft);
ASSERT_EQ(pair2.first, gold_value);
ASSERT_EQ(pair2.second, gold_location);
ASSERT_EQ(pair2.first, gold_value)
<< order_to_string(StdReducersTestEnumOrder::RightToLeft);
ASSERT_EQ(pair2.second, gold_location)
<< order_to_string(StdReducersTestEnumOrder::RightToLeft);
const auto pair3 = run_min_or_max_test<0, hostspace, index_type>(
view_h, StdReducersTestEnumOrder::Random);
ASSERT_EQ(pair3.first, gold_value);
ASSERT_EQ(pair3.second, gold_location);
ASSERT_EQ(pair3.first, gold_value)
<< order_to_string(StdReducersTestEnumOrder::Random);
ASSERT_EQ(pair3.second, gold_location)
<< order_to_string(StdReducersTestEnumOrder::Random);
}
TEST(std_algorithms_reducers, min_first_loc) {
@ -191,9 +194,6 @@ void run_min_max_test(ViewType view, StdReducersTestEnumOrder enValue,
static_assert(std::is_same<ExeSpace, Kokkos::HostSpace>::value,
"test is only enabled for HostSpace");
std::cout << "checking reduction with order: " << order_to_string(enValue)
<< "\n";
using view_value_type = typename ViewType::value_type;
using reducer_type =
Kokkos::MinMaxFirstLastLoc<view_value_type, IndexType, ExeSpace>;
@ -212,10 +212,10 @@ void run_min_max_test(ViewType view, StdReducersTestEnumOrder enValue,
reduction_value_type{view(index), view(index), index, index});
}
ASSERT_EQ(red_result.min_val, gold_values.first);
ASSERT_EQ(red_result.max_val, gold_values.second);
ASSERT_EQ(red_result.min_loc, gold_locs.first);
ASSERT_EQ(red_result.max_loc, gold_locs.second);
ASSERT_EQ(red_result.min_val, gold_values.first) << order_to_string(enValue);
ASSERT_EQ(red_result.max_val, gold_values.second) << order_to_string(enValue);
ASSERT_EQ(red_result.min_loc, gold_locs.first) << order_to_string(enValue);
ASSERT_EQ(red_result.max_loc, gold_locs.second) << order_to_string(enValue);
}
TEST(std_algorithms_reducers, min_max_first_last_loc) {