sorting an empty view should exit early and not fail (kokkos#6130)

* sorting an empty view should not fail

* fix binsort too

* fix unused

* fix test

* add comment

* address review comments

* add fences

* address comments

* address comments

* remove unused

algorithms/src/Kokkos_Sort.hpp

@@ -329,6 +329,10 @@ class BinSort {
   template <class ExecutionSpace, class ValuesViewType>
   void sort(const ExecutionSpace& exec, ValuesViewType const& values,
             int values_range_begin, int values_range_end) const {
+    if (values.extent(0) == 0) {
+      return;
+    }
+
     static_assert(
         Kokkos::SpaceAccessibility<ExecutionSpace,
                                    typename Space::memory_space>::accessible,
@@ -606,6 +610,10 @@ std::enable_if_t<(Kokkos::is_execution_space<ExecutionSpace>::value) &&
                                     memory_space>::accessible)>
 sort(const ExecutionSpace& exec,
      const Kokkos::View<DataType, Properties...>& view) {
+  if (view.extent(0) == 0) {
+    return;
+  }
+
   using ViewType = Kokkos::View<DataType, Properties...>;
   using CompType = BinOp1D<ViewType>;
 
@@ -648,8 +656,11 @@ sort(const ExecutionSpace& exec,
 template <class DataType, class... Properties>
 void sort(const Experimental::SYCL& space,
           const Kokkos::View<DataType, Properties...>& view) {
-  using ViewType = Kokkos::View<DataType, Properties...>;
+  if (view.extent(0) == 0) {
+    return;
+  }
 
+  using ViewType = Kokkos::View<DataType, Properties...>;
   static_assert(SpaceAccessibility<Experimental::SYCL,
                                    typename ViewType::memory_space>::accessible,
                 "SYCL execution space is not able to access the memory space "
@@ -676,6 +687,9 @@ std::enable_if_t<(Kokkos::is_execution_space<ExecutionSpace>::value) &&
                      HostSpace, typename Kokkos::View<DataType, Properties...>::
                                     memory_space>::accessible)>
 sort(const ExecutionSpace&, const Kokkos::View<DataType, Properties...>& view) {
+  if (view.extent(0) == 0) {
+    return;
+  }
   auto first = Experimental::begin(view);
   auto last  = Experimental::end(view);
   std::sort(first, last);
@@ -685,6 +699,9 @@ sort(const ExecutionSpace&, const Kokkos::View<DataType, Properties...>& view) {
 template <class DataType, class... Properties>
 void sort(const Cuda& space,
           const Kokkos::View<DataType, Properties...>& view) {
+  if (view.extent(0) == 0) {
+    return;
+  }
   const auto exec = thrust::cuda::par.on(space.cuda_stream());
   auto first      = Experimental::begin(view);
   auto last       = Experimental::end(view);
@@ -695,6 +712,11 @@ void sort(const Cuda& space,
 template <class ViewType>
 void sort(ViewType const& view) {
   Kokkos::fence("Kokkos::sort: before");
+
+  if (view.extent(0) == 0) {
+    return;
+  }
+
   typename ViewType::execution_space exec;
   sort(exec, view);
   exec.fence("Kokkos::sort: fence after sorting");
@@ -704,6 +726,10 @@ template <class ExecutionSpace, class ViewType>
 std::enable_if_t<Kokkos::is_execution_space<ExecutionSpace>::value> sort(
     const ExecutionSpace& exec, ViewType view, size_t const begin,
     size_t const end) {
+  if (view.extent(0) == 0) {
+    return;
+  }
+
   using range_policy = Kokkos::RangePolicy<typename ViewType::execution_space>;
   using CompType     = BinOp1D<ViewType>;
 
@@ -726,6 +752,11 @@ std::enable_if_t<Kokkos::is_execution_space<ExecutionSpace>::value> sort(
 template <class ViewType>
 void sort(ViewType view, size_t const begin, size_t const end) {
   Kokkos::fence("Kokkos::sort: before");
+
+  if (view.extent(0) == 0) {
+    return;
+  }
+
   typename ViewType::execution_space exec;
   sort(exec, view, begin, end);
   exec.fence("Kokkos::Sort: fence after sorting");

algorithms/unit_tests/TestSort.hpp

@@ -425,6 +425,51 @@ void test_sort(unsigned int N) {
   test_sort_integer_overflow<ExecutionSpace, unsigned long long>();
   test_sort_integer_overflow<ExecutionSpace, int>();
 }
+
+template <class ExecutionSpace>
+void test_sort_empty_view() {
+  // does not matter if we use int or something else
+  Kokkos::View<int*, ExecutionSpace> v("v", 0);
+
+  // TODO check the synchronous behavior of the calls below
+  ASSERT_NO_THROW(Kokkos::sort(ExecutionSpace(), v));
+  ASSERT_NO_THROW(Kokkos::sort(v));
+}
+
+template <class ExecutionSpace>
+void test_binsort_empty_view() {
+  // the bounds and extents used below are totally arbitrary
+  // and, in theory, should have no impact
+
+  using KeyViewType = Kokkos::View<int*, ExecutionSpace>;
+  KeyViewType kv("kv", 20);
+
+  using BinOp_t = Kokkos::BinOp1D<KeyViewType>;
+  BinOp_t binOp(5, 0, 10);
+  Kokkos::BinSort<KeyViewType, BinOp_t> Sorter(ExecutionSpace{}, kv, binOp);
+
+  // does not matter if we use int or something else
+  Kokkos::View<int*, ExecutionSpace> v("v", 0);
+
+  // test all exposed public sort methods
+  ASSERT_NO_THROW(Sorter.sort(ExecutionSpace(), v, 0, 0));
+  ASSERT_NO_THROW(Sorter.sort(v, 0, 0));
+  ASSERT_NO_THROW(Sorter.sort(ExecutionSpace(), v));
+  ASSERT_NO_THROW(Sorter.sort(v));
+}
+
+template <class ExecutionSpace>
+void test_binsort_empty_keys() {
+  using KeyViewType = Kokkos::View<int*, ExecutionSpace>;
+  KeyViewType kv("kv", 0);
+
+  using BinOp_t = Kokkos::BinOp1D<KeyViewType>;
+  BinOp_t binOp(5, 0, 10);
+  Kokkos::BinSort<KeyViewType, BinOp_t> Sorter(ExecutionSpace{}, kv, binOp);
+
+  ASSERT_NO_THROW(Sorter.create_permute_vector(ExecutionSpace{}));
+}
+
 }  // namespace Impl
 }  // namespace Test
 #endif /* KOKKOS_ALGORITHMS_UNITTESTS_TESTSORT_HPP */

algorithms/unit_tests/TestSortCommon.hpp

@@ -23,5 +23,17 @@ namespace Test {
 TEST(TEST_CATEGORY, SortUnsigned) {
   Impl::test_sort<TEST_EXECSPACE, unsigned>(171);
 }
+
+TEST(TEST_CATEGORY, SortEmptyView) {
+  Impl::test_sort_empty_view<TEST_EXECSPACE>();
+}
+
+TEST(TEST_CATEGORY, BinSortEmptyView) {
+  Impl::test_binsort_empty_view<TEST_EXECSPACE>();
+}
+
+TEST(TEST_CATEGORY, BinSortEmptyKeys) {
+  Impl::test_binsort_empty_keys<TEST_EXECSPACE>();
+}
 }  // namespace Test
 #endif