Asynchronous algorithms testing enhancements:
* Rename `DECLARE_VARIABLE_UNITTEST_WITH_TYPES_AND_NAME`
to `DECLARE_GENERIC_UNITTEST_WITH_TYPES_AND_NAME`.
* Add `DECLARE_GENERIC_SIZED_UNITTEST_WITH_TYPES_AND_NAME`.
* Add `thrust::async::transform` tests.
Bug
Bug
MODIFY
testing/async_copy.cu
testing/async_copy.cu
@@ -77,12 +77,12 @@ struct test_async_copy_host_to_device
}
};
};
-DECLARE_VARIABLE_UNITTEST_WITH_TYPES_AND_NAME(
+DECLARE_GENERIC_SIZED_UNITTEST_WITH_TYPES_AND_NAME(
test_async_copy_host_to_device<invoke_async_copy_fn>::tester
, BuiltinNumericTypes
, test_async_copy_trivially_relocatable_elements_host_to_device
);
-DECLARE_VARIABLE_UNITTEST_WITH_TYPES_AND_NAME(
+DECLARE_GENERIC_SIZED_UNITTEST_WITH_TYPES_AND_NAME(
test_async_copy_host_to_device<invoke_async_copy_host_to_device_fn>::tester
, BuiltinNumericTypes
, test_async_copy_trivially_relocatable_elements_host_to_device_policies
@@ -118,12 +118,12 @@ struct test_async_copy_device_to_host
}
};
};
-DECLARE_VARIABLE_UNITTEST_WITH_TYPES_AND_NAME(
+DECLARE_GENERIC_SIZED_UNITTEST_WITH_TYPES_AND_NAME(
test_async_copy_device_to_host<invoke_async_copy_fn>::tester
, BuiltinNumericTypes
, test_async_copy_trivially_relocatable_elements_device_to_host
);
-DECLARE_VARIABLE_UNITTEST_WITH_TYPES_AND_NAME(
+DECLARE_GENERIC_SIZED_UNITTEST_WITH_TYPES_AND_NAME(
test_async_copy_device_to_host<invoke_async_copy_device_to_host_fn>::tester
, BuiltinNumericTypes
, test_async_copy_trivially_relocatable_elements_device_to_host_policies
@@ -159,17 +159,17 @@ struct test_async_copy_device_to_device
}
};
};
-DECLARE_VARIABLE_UNITTEST_WITH_TYPES_AND_NAME(
+DECLARE_GENERIC_SIZED_UNITTEST_WITH_TYPES_AND_NAME(
test_async_copy_device_to_device<invoke_async_copy_fn>::tester
, NumericTypes
, test_async_copy_device_to_device
);
-DECLARE_VARIABLE_UNITTEST_WITH_TYPES_AND_NAME(
+DECLARE_GENERIC_SIZED_UNITTEST_WITH_TYPES_AND_NAME(
test_async_copy_device_to_device<invoke_async_copy_device_fn>::tester
, NumericTypes
, test_async_copy_device_to_device_policy
);
-DECLARE_VARIABLE_UNITTEST_WITH_TYPES_AND_NAME(
+DECLARE_GENERIC_SIZED_UNITTEST_WITH_TYPES_AND_NAME(
test_async_copy_device_to_device<invoke_async_copy_device_to_device_fn>::tester
, NumericTypes
, test_async_copy_device_to_device_policies
@@ -207,28 +207,28 @@ struct test_async_copy_counting_iterator_input_to_device_vector
}
};
};
-DECLARE_VARIABLE_UNITTEST_WITH_TYPES_AND_NAME(
+DECLARE_GENERIC_SIZED_UNITTEST_WITH_TYPES_AND_NAME(
test_async_copy_counting_iterator_input_to_device_vector<
invoke_async_copy_fn
>::tester
, BuiltinNumericTypes
, test_async_copy_counting_iterator_input_trivially_relocatable_elements_device_to_device
);
-DECLARE_VARIABLE_UNITTEST_WITH_TYPES_AND_NAME(
+DECLARE_GENERIC_SIZED_UNITTEST_WITH_TYPES_AND_NAME(
test_async_copy_counting_iterator_input_to_device_vector<
invoke_async_copy_device_fn
>::tester
, BuiltinNumericTypes
, test_async_copy_counting_iterator_input_trivially_relocatable_elements_device_to_device_policy
);
-DECLARE_VARIABLE_UNITTEST_WITH_TYPES_AND_NAME(
+DECLARE_GENERIC_SIZED_UNITTEST_WITH_TYPES_AND_NAME(
test_async_copy_counting_iterator_input_to_device_vector<
invoke_async_copy_device_to_device_fn
>::tester
, BuiltinNumericTypes
, test_async_copy_counting_iterator_input_trivially_relocatable_elements_device_to_device_policies
);
-DECLARE_VARIABLE_UNITTEST_WITH_TYPES_AND_NAME(
+DECLARE_GENERIC_SIZED_UNITTEST_WITH_TYPES_AND_NAME(
test_async_copy_counting_iterator_input_to_device_vector<
invoke_async_copy_host_to_device_fn
>::tester
@@ -269,14 +269,14 @@ struct test_async_copy_counting_iterator_input_to_host_vector
}
};
};
-DECLARE_VARIABLE_UNITTEST_WITH_TYPES_AND_NAME(
+DECLARE_GENERIC_SIZED_UNITTEST_WITH_TYPES_AND_NAME(
test_async_copy_counting_iterator_input_to_host_vector<
invoke_async_copy_fn
>::tester
, BuiltinNumericTypes
, test_async_copy_counting_iterator_input_trivially_relocatable_elements_device_to_host
);
-DECLARE_VARIABLE_UNITTEST_WITH_TYPES_AND_NAME(
+DECLARE_GENERIC_SIZED_UNITTEST_WITH_TYPES_AND_NAME(
test_async_copy_counting_iterator_input_to_host_vector<
invoke_async_copy_device_to_host_fn
>::tester

MODIFY
testing/async_for_each.cu
testing/async_for_each.cu
@@ -38,7 +38,7 @@ DEFINE_ASYNC_FOR_EACH_CALLABLE(
///////////////////////////////////////////////////////////////////////////////
-struct divide_by_2
+struct inplace_divide_by_2
{
template <typename T>
__host__ __device__
@@ -74,18 +74,24 @@ struct test_async_for_each
}
};
};
-DECLARE_VARIABLE_UNITTEST_WITH_TYPES_AND_NAME(
+DECLARE_GENERIC_SIZED_UNITTEST_WITH_TYPES_AND_NAME(
THRUST_PP_EXPAND_ARGS(
- test_async_for_each<invoke_async_for_each_fn, divide_by_2>::tester
+ test_async_for_each<
+ invoke_async_for_each_fn
+ , inplace_divide_by_2
+ >::tester
)
-, BuiltinNumericTypes
+, NumericTypes
, test_async_for_each
);
-DECLARE_VARIABLE_UNITTEST_WITH_TYPES_AND_NAME(
+DECLARE_GENERIC_SIZED_UNITTEST_WITH_TYPES_AND_NAME(
THRUST_PP_EXPAND_ARGS(
- test_async_for_each<invoke_async_for_each_device_fn, divide_by_2>::tester
+ test_async_for_each<
+ invoke_async_for_each_device_fn
+ , inplace_divide_by_2
+ >::tester
)
-, BuiltinNumericTypes
+, NumericTypes
, test_async_for_each_policy
);

ADD
testing/async_transform.cu
+#include <thrust/detail/config.h>
+
+#if THRUST_CPP_DIALECT >= 2011
+
+#include <unittest/unittest.h>
+
+#include <thrust/async/transform.h>
+#include <thrust/host_vector.h>
+#include <thrust/device_vector.h>
+
+#define DEFINE_ASYNC_TRANSFORM_CALLABLE(name, ...) \
+ struct THRUST_PP_CAT2(name, _fn) \
+ { \
+ template < \
+ typename ForwardIt, typename Sentinel, typename OutputIt \
+ , typename UnaryOperation \
+ > \
+ __host__ \
+ auto operator()( \
+ ForwardIt&& first, Sentinel&& last, OutputIt&& output \
+ , UnaryOperation&& f \
+ ) const \
+ THRUST_DECLTYPE_RETURNS( \
+ ::thrust::async::transform( \
+ __VA_ARGS__ \
+ THRUST_PP_COMMA_IF(THRUST_PP_ARITY(__VA_ARGS__)) \
+ THRUST_FWD(first), THRUST_FWD(last), THRUST_FWD(output), THRUST_FWD(f)\
+ ) \
+ ) \
+ }; \
+ /**/
+
+DEFINE_ASYNC_TRANSFORM_CALLABLE(
+ invoke_async_transform
+);
+
+DEFINE_ASYNC_TRANSFORM_CALLABLE(
+ invoke_async_transform_device, thrust::device
+);
+
+#undef DEFINE_ASYNC_TRANSFORM_CALLABLE
+
+///////////////////////////////////////////////////////////////////////////////
+
+struct divide_by_2
+{
+ template <typename T>
+ __host__ __device__
+ T operator()(T x) const
+ {
+ return x / 2;
+ }
+};
+
+///////////////////////////////////////////////////////////////////////////////
+
+template <typename AsyncTransformCallable, typename UnaryOperation>
+struct test_async_transform_unary
+{
+ 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);
+
+ thrust::host_vector<T> h1_data(n);
+ thrust::device_vector<T> d1_data(n);
+
+ thrust::transform(
+ h0_data.begin(), h0_data.end(), h1_data.begin(), UnaryOperation{}
+ );
+
+ auto f0 = AsyncTransformCallable{}(
+ d0_data.begin(), d0_data.end(), d1_data.begin(), UnaryOperation{}
+ );
+
+ f0.wait();
+
+ ASSERT_EQUAL(h0_data, d0_data);
+ ASSERT_EQUAL(h1_data, d1_data);
+ }
+ };
+};
+DECLARE_GENERIC_SIZED_UNITTEST_WITH_TYPES_AND_NAME(
+ THRUST_PP_EXPAND_ARGS(
+ test_async_transform_unary<
+ invoke_async_transform_fn
+ , divide_by_2
+ >::tester
+ )
+, NumericTypes
+, test_async_transform_unary
+);
+DECLARE_GENERIC_SIZED_UNITTEST_WITH_TYPES_AND_NAME(
+ THRUST_PP_EXPAND_ARGS(
+ test_async_transform_unary<
+ invoke_async_transform_device_fn
+ , divide_by_2
+ >::tester
+ )
+, NumericTypes
+, test_async_transform_unary_policy
+);
+
+///////////////////////////////////////////////////////////////////////////////
+
+template <typename AsyncTransformCallable, typename UnaryOperation>
+struct test_async_transform_unary_inplace
+{
+ 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);
+
+ thrust::transform(
+ h0_data.begin(), h0_data.end(), h0_data.begin(), UnaryOperation{}
+ );
+
+ auto f0 = AsyncTransformCallable{}(
+ d0_data.begin(), d0_data.end(), d0_data.begin(), UnaryOperation{}
+ );
+
+ f0.wait();
+
+ ASSERT_EQUAL(h0_data, d0_data);
+ }
+ };
+};
+DECLARE_GENERIC_SIZED_UNITTEST_WITH_TYPES_AND_NAME(
+ THRUST_PP_EXPAND_ARGS(
+ test_async_transform_unary_inplace<
+ invoke_async_transform_fn
+ , divide_by_2
+ >::tester
+ )
+, NumericTypes
+, test_async_transform_unary_inplace
+);
+DECLARE_GENERIC_SIZED_UNITTEST_WITH_TYPES_AND_NAME(
+ THRUST_PP_EXPAND_ARGS(
+ test_async_transform_unary_inplace<
+ invoke_async_transform_device_fn
+ , divide_by_2
+ >::tester
+ )
+, NumericTypes
+, test_async_transform_unary_inplace_policy
+);
+
+#endif // THRUST_CPP_DIALECT >= 2011
+

MODIFY
testing/unittest/testframework.h
testing/unittest/testframework.h
@@ -478,7 +478,11 @@ class TEST##UnitTest : public UnitTest { \
}; \
TEST##UnitTest TEST##Instance
-#define DECLARE_VARIABLE_UNITTEST_WITH_TYPES_AND_NAME(TEST, TYPES, NAME) \
+#define DECLARE_GENERIC_UNITTEST_WITH_TYPES_AND_NAME(TEST, TYPES, NAME) \
+ ::SimpleUnitTest<TEST, TYPES> NAME##_instance(#NAME) \
+ /**/
+
+#define DECLARE_GENERIC_SIZED_UNITTEST_WITH_TYPES_AND_NAME(TEST, TYPES, NAME) \
::VariableUnitTest<TEST, TYPES> NAME##_instance(#NAME) \
/**/

Asynchronous algorithms testing enhancements:

* Rename `DECLARE_VARIABLE_UNITTEST_WITH_TYPES_AND_NAME`
  to `DECLARE_GENERIC_UNITTEST_WITH_TYPES_AND_NAME`.
* Add `DECLARE_GENERIC_SIZED_UNITTEST_WITH_TYPES_AND_NAME`.
* Add `thrust::async::transform` tests.

Bug 2379510
Bug 2463967