Add initial unit tests for `thrust::async::sort`.
Bug
ADD
testing/async_reduce.cu
+#include <thrust/detail/config.h>
+
+#if THRUST_CPP_DIALECT >= 2011
+
+#include <unittest/unittest.h>
+
+#include <thrust/async/reduce.h>
+#include <thrust/host_vector.h>
+#include <thrust/device_vector.h>
+
+template <typename T>
+struct custom_plus
+{
+ __host__ __device__
+ T operator()(T lhs, T rhs) const
+ {
+ return lhs + rhs;
+ }
+};
+
+template <typename T>
+struct test_async_reduce
+{
+ __host__
+ void operator()(std::size_t n)
+ {
+ thrust::host_vector<T> h0_data(unittest::random_integers<T>(n));
+ thrust::device_vector<T> d0_data(h0_data);
+
+ ASSERT_EQUAL(h0_data, d0_data);
+
+ auto r0 = thrust::reduce(
+ h0_data.begin(), h0_data.end()
+ );
+
+ auto f0 = thrust::async::reduce(
+ d0_data.begin(), d0_data.end()
+ );
+
+ auto r1 = std::move(f0).get();
+
+ ASSERT_EQUAL(r0, r1);
+ }
+};
+// TODO: Switch to `DECLARE_VARIABLE_UNITTEST` when we add `custom_numeric` to
+// the list of types it covers.
+VariableUnitTest<
+ test_async_reduce
+, NumericTypes
+> test_async_reduce_instance;
+
+template <typename T>
+struct test_async_reduce_with_policy
+{
+ __host__
+ void operator()(std::size_t n)
+ {
+ thrust::host_vector<T> h0_data(unittest::random_integers<T>(n));
+ thrust::device_vector<T> d0_data(h0_data);
+
+ ASSERT_EQUAL(h0_data, d0_data);
+
+ auto r0 = thrust::reduce(
+ h0_data.begin(), h0_data.end()
+ );
+
+ auto f0 = thrust::async::reduce(
+ thrust::device, d0_data.begin(), d0_data.end()
+ );
+
+ auto r1 = std::move(f0).get();
+
+ ASSERT_EQUAL(r0, r1);
+ }
+};
+// TODO: Switch to `DECLARE_VARIABLE_UNITTEST` when we add `custom_numeric` to
+// the list of types it covers.
+VariableUnitTest<
+ test_async_reduce_with_policy
+, NumericTypes
+> test_async_reduce_with_policy_instance;
+
+template <typename T>
+struct test_async_reduce_with_init
+{
+ __host__
+ void operator()(std::size_t n)
+ {
+ thrust::host_vector<T> h0_data(unittest::random_integers<T>(n));
+ thrust::device_vector<T> d0_data(h0_data);
+
+ ASSERT_EQUAL(h0_data, d0_data);
+
+ T const init = unittest::random_integer<T>();
+
+ auto r0 = thrust::reduce(
+ h0_data.begin(), h0_data.end(), init
+ );
+
+ auto f0 = thrust::async::reduce(
+ d0_data.begin(), d0_data.end(), init
+ );
+
+ auto r1 = std::move(f0).get();
+
+ ASSERT_EQUAL(r0, r1);
+ }
+};
+// TODO: Switch to `DECLARE_VARIABLE_UNITTEST` when we add `custom_numeric` to
+// the list of types it covers.
+VariableUnitTest<
+ test_async_reduce_with_init
+, NumericTypes
+> test_async_reduce_with_init_instance;
+
+template <typename T>
+struct test_async_reduce_with_policy_init
+{
+ __host__
+ void operator()(std::size_t n)
+ {
+ thrust::host_vector<T> h0_data(unittest::random_integers<T>(n));
+ thrust::device_vector<T> d0_data(h0_data);
+
+ ASSERT_EQUAL(h0_data, d0_data);
+
+ T const init = unittest::random_integer<T>();
+
+ auto r0 = thrust::reduce(
+ h0_data.begin(), h0_data.end(), init
+ );
+
+ auto f0 = thrust::async::reduce(
+ thrust::device, d0_data.begin(), d0_data.end(), init
+ );
+
+ auto r1 = std::move(f0).get();
+
+ ASSERT_EQUAL(r0, r1);
+ }
+};
+// TODO: Switch to `DECLARE_VARIABLE_UNITTEST` when we add `custom_numeric` to
+// the list of types it covers.
+VariableUnitTest<
+ test_async_reduce_with_policy_init
+, NumericTypes
+> test_async_reduce_with_policy_init_instance;
+
+template <typename T>
+struct test_async_reduce_with_init_op
+{
+ __host__
+ void operator()(std::size_t n)
+ {
+ thrust::host_vector<T> h0_data(unittest::random_integers<T>(n));
+ thrust::device_vector<T> d0_data(h0_data);
+
+ ASSERT_EQUAL(h0_data, d0_data);
+
+ T const init = unittest::random_integer<T>();
+ custom_plus<T> op{};
+
+ auto r0 = thrust::reduce(
+ h0_data.begin(), h0_data.end(), init, op
+ );
+
+ auto f0 = thrust::async::reduce(
+ d0_data.begin(), d0_data.end(), init, op
+ );
+
+ auto r1 = std::move(f0).get();
+
+ ASSERT_EQUAL(r0, r1);
+ }
+};
+// TODO: Switch to `DECLARE_VARIABLE_UNITTEST` when we add `custom_numeric` to
+// the list of types it covers.
+VariableUnitTest<
+ test_async_reduce_with_init_op
+, NumericTypes
+> test_async_reduce_with_init_op_instance;
+
+template <typename T>
+struct test_async_reduce_with_policy_init_op
+{
+ __host__
+ void operator()(std::size_t n)
+ {
+ thrust::host_vector<T> h0_data(unittest::random_integers<T>(n));
+ thrust::device_vector<T> d0_data(h0_data);
+
+ ASSERT_EQUAL(h0_data, d0_data);
+
+ T const init = unittest::random_integer<T>();
+ custom_plus<T> op{};
+
+ auto r0 = thrust::reduce(
+ h0_data.begin(), h0_data.end(), init, op
+ );
+
+ auto f0 = thrust::async::reduce(
+ thrust::device, d0_data.begin(), d0_data.end(), init, op
+ );
+
+ auto r1 = std::move(f0).get();
+
+ ASSERT_EQUAL(r0, r1);
+ }
+};
+// TODO: Switch to `DECLARE_VARIABLE_UNITTEST` when we add `custom_numeric` to
+// the list of types it covers.
+VariableUnitTest<
+ test_async_reduce_with_policy_init_op
+, NumericTypes
+> test_async_reduce_with_policy_init_op_instance;
+
+// TODO: Async copy then reduce.
+
+// TODO: Device-side reduction usage.
+
+// TODO: Make random_integers more generic.
+
+#endif // THRUST_CPP_DIALECT >= 2011
+

ADD
testing/async_sort.cu
+#include <thrust/detail/config.h>
+
+#if THRUST_CPP_DIALECT >= 2011
+
+#include <unittest/unittest.h>
+
+#include <thrust/async/sort.h>
+#include <thrust/host_vector.h>
+#include <thrust/device_vector.h>
+
+template <typename T>
+struct custom_greater
+{
+ __host__ __device__
+ bool operator()(T rhs, T lhs) const
+ {
+ return lhs > rhs;
+ }
+};
+
+template <typename T>
+struct test_async_sort
+{
+ __host__
+ void operator()(std::size_t n)
+ {
+ thrust::host_vector<T> h0_data(unittest::random_integers<T>(n));
+ thrust::device_vector<T> d0_data(h0_data);
+
+ ASSERT_EQUAL(h0_data, d0_data);
+
+ thrust::sort(
+ h0_data.begin(), h0_data.end()
+ );
+
+ auto f0 = thrust::async::sort(
+ d0_data.begin(), d0_data.end()
+ );
+
+ f0.wait();
+
+ ASSERT_EQUAL(h0_data, d0_data);
+ }
+};
+// TODO: Switch to `DECLARE_VARIABLE_UNITTEST` when we add `custom_numeric` to
+// the list of types it covers.
+VariableUnitTest<
+ test_async_sort
+, NumericTypes
+> test_async_sort_instance;
+
+template <typename T>
+struct test_async_sort_with_policy
+{
+ __host__
+ void operator()(std::size_t n)
+ {
+ thrust::host_vector<T> h0_data(unittest::random_integers<T>(n));
+ thrust::device_vector<T> d0_data(h0_data);
+
+ ASSERT_EQUAL(h0_data, d0_data);
+
+ thrust::sort(
+ h0_data.begin(), h0_data.end()
+ );
+
+ auto f0 = thrust::async::sort(
+ thrust::device, d0_data.begin(), d0_data.end()
+ );
+
+ f0.wait();
+
+ ASSERT_EQUAL(h0_data, d0_data);
+ }
+};
+// TODO: Switch to `DECLARE_VARIABLE_UNITTEST` when we add `custom_numeric` to
+// the list of types it covers.
+VariableUnitTest<
+ test_async_sort_with_policy
+, NumericTypes
+> test_async_sort_with_policy_instance;
+
+template <template <typename> class Op>
+struct test_async_sort_with_op
+{
+ template <typename T>
+ struct tester
+ {
+ __host__
+ void operator()(std::size_t n)
+ {
+ thrust::host_vector<T> h0_data(unittest::random_integers<T>(n));
+ thrust::device_vector<T> d0_data(h0_data);
+
+ ASSERT_EQUAL(h0_data, d0_data);
+
+ Op<T> op{};
+
+ thrust::sort(
+ h0_data.begin(), h0_data.end(), op
+ );
+
+ auto f0 = thrust::async::sort(
+ d0_data.begin(), d0_data.end(), op
+ );
+
+ f0.wait();
+
+ ASSERT_EQUAL(h0_data, d0_data);
+ }
+ };
+};
+// TODO: Switch to `DECLARE_VARIABLE_UNITTEST` when we add `custom_numeric` to
+// the list of types it covers.
+VariableUnitTest<
+ test_async_sort_with_op<custom_greater>::tester
+, NumericTypes
+> test_async_sort_with_op_instance(
+ "test_async_sort_with_op<custom_greater>"
+);
+VariableUnitTest<
+ test_async_sort_with_op<thrust::less>::tester
+, NumericTypes
+> test_async_sort_with_less_instance(
+ "test_async_sort_with_op<thrust::less>"
+);
+VariableUnitTest<
+ test_async_sort_with_op<thrust::greater>::tester
+, NumericTypes
+> test_async_sort_with_greater_instance(
+ "test_async_sort_with_op<thrust::greater>"
+);
+
+template <template <typename> class Op>
+struct test_async_sort_with_policy_op
+{
+ template <typename T>
+ struct tester
+ {
+ __host__
+ void operator()(std::size_t n)
+ {
+ thrust::host_vector<T> h0_data(unittest::random_integers<T>(n));
+ thrust::device_vector<T> d0_data(h0_data);
+
+ ASSERT_EQUAL(h0_data, d0_data);
+
+ Op<T> op{};
+
+ thrust::sort(
+ h0_data.begin(), h0_data.end(), op
+ );
+
+ auto f0 = thrust::async::sort(
+ thrust::device, d0_data.begin(), d0_data.end(), op
+ );
+
+ f0.wait();
+
+ ASSERT_EQUAL(h0_data, d0_data);
+ }
+ };
+};
+// TODO: Switch to `DECLARE_VARIABLE_UNITTEST` when we add `custom_numeric` to
+// the list of types it covers.
+VariableUnitTest<
+ test_async_sort_with_policy_op<custom_greater>::tester
+, NumericTypes
+> test_async_sort_with_policy_op_instance(
+ "test_async_sort_with_policy_op<custom_greater>"
+);
+VariableUnitTest<
+ test_async_sort_with_policy_op<thrust::less>::tester
+, NumericTypes
+> test_async_sort_with_policy_less_instance(
+ "test_async_sort_with_policy_op<thrust::less>"
+);
+VariableUnitTest<
+ test_async_sort_with_policy_op<thrust::greater>::tester
+, NumericTypes
+> test_async_sort_with_policy_greater_instance(
+ "test_async_sort_with_policy_op<thrust::greater>"
+);
+
+// TODO: Async copy then sort.
+
+// TODO: Test future return type.
+
+#endif // THRUST_CPP_DIALECT >= 2011
+

MODIFY
testing/reduce.cu
testing/reduce.cu
@@ -8,7 +8,7 @@ template<typename T>
struct plus_mod_10
{
__host__ __device__
- T operator()(T rhs, T lhs) const
+ T operator()(T lhs, T rhs) const
{
return ((lhs % 10) + (rhs % 10)) % 10;
}

MODIFY
testing/unittest/testframework.h
testing/unittest/testframework.h
@@ -474,6 +474,9 @@ template<template <typename> class TestName, typename TypeList>
VariableUnitTest()
: UnitTest(base_class_name(unittest::type_name<TestName<int> >()).c_str()) {}
+ VariableUnitTest(const char * name)
+ : UnitTest(name) {}
+
void run()
{
std::vector<size_t> sizes = get_test_sizes();

MODIFY
thrust/async/reduce.h
thrust/async/reduce.h
@@ -161,9 +161,8 @@ struct reduce_fn final
template <typename ForwardIt, typename Sentinel>
__host__ __device__
static auto call(ForwardIt&& first, Sentinel&& last)
- THRUST_DECLTYPE_RETURNS_WITH_SFINAE_CONDITION(
- (negation<is_execution_policy<remove_cvref_t<ForwardIt>>>::value)
- , call(
+ THRUST_DECLTYPE_RETURNS(
+ call(
THRUST_FWD(first), THRUST_FWD(last)
, typename iterator_traits<remove_cvref_t<ForwardIt>>::value_type{}
)

MODIFY
thrust/async/sort.h
thrust/async/sort.h
#include <thrust/detail/static_assert.h>
#include <thrust/detail/select_system.h>
+#include <thrust/type_traits/logical_metafunctions.h>
+#include <thrust/type_traits/remove_cvref.h>
+#include <thrust/type_traits/is_execution_policy.h>
#include <thrust/system/detail/adl/async/sort.h>
#include <thrust/future.h>
@@ -206,8 +209,9 @@ struct sort_fn final
template <typename ForwardIt, typename Sentinel, typename StrictWeakOrdering>
__host__ __device__
static auto call(ForwardIt&& first, Sentinel&& last, StrictWeakOrdering&& comp)
- THRUST_DECLTYPE_RETURNS(
- call(
+ THRUST_DECLTYPE_RETURNS_WITH_SFINAE_CONDITION(
+ (negation<is_execution_policy<remove_cvref_t<ForwardIt>>>::value)
+ , call(
thrust::detail::select_system(
typename thrust::iterator_system<ForwardIt>::type{}
)

MODIFY
thrust/system/cuda/detail/async/sort.h
thrust/system/cuda/detail/async/sort.h
@@ -102,7 +102,7 @@ auto async_stable_sort_n(
Size n,
StrictWeakOrdering comp
) ->
- typename std::enable_if<
+/* typename std::enable_if<
conjunction<
is_contiguous_iterator<ForwardIt>
, negation<
@@ -111,13 +111,15 @@ auto async_stable_sort_n(
>
>
>::value
- , unique_eager_future<
+ ,
+*/
+ unique_eager_future<
void
, typename thrust::detail::allocator_traits<
decltype(get_async_device_allocator(policy))
>::template rebind_traits<void>::pointer
>
- >::type
+// >::type
{
using T = typename thrust::iterator_traits<ForwardIt>::value_type;

Add initial unit tests for `thrust::async::sort`.

Bug 2379510