Added transform async tests.
MODIFY
test/test_async_copy.cpp
test/test_async_copy.cpp
@@ -84,12 +84,14 @@ void AsyncCopyHostToDevice()
TYPED_TEST(AsyncCopyTests, TestAsyncTriviallyRelocatableElementsHostToDevice)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
AsyncCopyHostToDevice<T, invoke_async_copy_fn>();
};
TYPED_TEST(AsyncCopyTests, TestAsyncTriviallyRelocatableElementsHostToDevicePolicies)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
AsyncCopyHostToDevice<T, invoke_async_copy_host_to_device_fn>();
};
@@ -129,12 +131,14 @@ void AsyncCopyDeviceToHost()
TYPED_TEST(AsyncCopyTests, TestAsyncCopyTriviallyRelocatableDeviceToHost)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
AsyncCopyDeviceToHost<T, invoke_async_copy_fn>();
};
TYPED_TEST(AsyncCopyTests, TestAsyncCopyTriviallyRelocatableDeviceToHostPolicies)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
AsyncCopyDeviceToHost<T, invoke_async_copy_device_to_host_fn>();
};

MODIFY
test/test_async_for_each.cpp
test/test_async_for_each.cpp
@@ -70,16 +70,14 @@ void test_async_for_each()
f0.wait();
- for(size_t i = 0; i < size; i++)
- {
- ASSERT_EQ(h0_data[i], d0_data[i]);
- }
+ ASSERT_EQ(h0_data, d0_data);
}
}
};
TYPED_TEST(AsyncForEachTests, TestAsyncForEach)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
test_async_for_each<
T,
@@ -90,6 +88,7 @@ TYPED_TEST(AsyncForEachTests, TestAsyncForEach)
TYPED_TEST(AsyncForEachTests, TestAsyncForEachPolicy)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
test_async_for_each<
T,
@@ -99,4 +98,3 @@ TYPED_TEST(AsyncForEachTests, TestAsyncForEachPolicy)
}
#endif
-

MODIFY
test/test_async_reduce.cpp
test/test_async_reduce.cpp
@@ -407,6 +407,7 @@ void testAsyncReduce()
TYPED_TEST(AsyncReduceTests, TestAsyncReduce)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
testAsyncReduce<
T,
@@ -417,6 +418,7 @@ TYPED_TEST(AsyncReduceTests, TestAsyncReduce)
TYPED_TEST(AsyncReduceTests, TestAsyncReducePolicy)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
testAsyncReduce<
T,
@@ -427,6 +429,7 @@ TYPED_TEST(AsyncReduceTests, TestAsyncReducePolicy)
TYPED_TEST(AsyncReduceTests, TestAsyncReducePolicyAllocator)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
testAsyncReduce<
T,
@@ -437,6 +440,7 @@ TYPED_TEST(AsyncReduceTests, TestAsyncReducePolicyAllocator)
TYPED_TEST(AsyncReduceTests, TestAsyncReducePolicyOn)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
testAsyncReduce<
T,
@@ -447,6 +451,7 @@ TYPED_TEST(AsyncReduceTests, TestAsyncReducePolicyOn)
TYPED_TEST(AsyncReduceTests, TestAsyncReducePolicyAllocatorOn)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
testAsyncReduce<
T,
@@ -457,6 +462,7 @@ TYPED_TEST(AsyncReduceTests, TestAsyncReducePolicyAllocatorOn)
TYPED_TEST(AsyncReduceTests, TestAsyncReduceInit)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
testAsyncReduce<
T,
@@ -467,6 +473,7 @@ TYPED_TEST(AsyncReduceTests, TestAsyncReduceInit)
TYPED_TEST(AsyncReduceTests, TestAsyncReducePolicyInit)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
testAsyncReduce<
T,
@@ -477,6 +484,7 @@ TYPED_TEST(AsyncReduceTests, TestAsyncReducePolicyInit)
TYPED_TEST(AsyncReduceTests, TestAsyncReducePolicyAllocatorInit)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
testAsyncReduce<
T,
@@ -487,6 +495,7 @@ TYPED_TEST(AsyncReduceTests, TestAsyncReducePolicyAllocatorInit)
TYPED_TEST(AsyncReduceTests, TestAsyncReducePolicyOnInit)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
testAsyncReduce<
T,
@@ -497,6 +506,7 @@ TYPED_TEST(AsyncReduceTests, TestAsyncReducePolicyOnInit)
TYPED_TEST(AsyncReduceTests, TestAsyncReducePolicyAllocatorOnInit)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
testAsyncReduce<
T,
@@ -507,6 +517,7 @@ TYPED_TEST(AsyncReduceTests, TestAsyncReducePolicyAllocatorOnInit)
TYPED_TEST(AsyncReduceTests, TestAsyncReduceInitPlus)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
testAsyncReduce<
T,
@@ -517,6 +528,7 @@ TYPED_TEST(AsyncReduceTests, TestAsyncReduceInitPlus)
TYPED_TEST(AsyncReduceTests, TestAsyncReducePolicyInitPlus)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
testAsyncReduce<
T,
@@ -527,6 +539,7 @@ TYPED_TEST(AsyncReduceTests, TestAsyncReducePolicyInitPlus)
TYPED_TEST(AsyncReduceTests, TestAsyncReducePolicyAllocatorInitPlus)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
testAsyncReduce<
T,
@@ -537,6 +550,7 @@ TYPED_TEST(AsyncReduceTests, TestAsyncReducePolicyAllocatorInitPlus)
TYPED_TEST(AsyncReduceTests, TestAsyncReducePolicyOnInitPlus)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
testAsyncReduce<
T,
@@ -547,6 +561,7 @@ TYPED_TEST(AsyncReduceTests, TestAsyncReducePolicyOnInitPlus)
TYPED_TEST(AsyncReduceTests, TestAsyncReducePolicyAllocatorOnInitPlus)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
testAsyncReduce<
T,
@@ -557,6 +572,7 @@ TYPED_TEST(AsyncReduceTests, TestAsyncReducePolicyAllocatorOnInitPlus)
TYPED_TEST(AsyncReduceTests, TestAsyncReduceInitCustomPlus)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
testAsyncReduce<
T,
@@ -567,6 +583,7 @@ TYPED_TEST(AsyncReduceTests, TestAsyncReduceInitCustomPlus)
TYPED_TEST(AsyncReduceTests, TestAsyncReducePolicyInitCustomPlus)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
testAsyncReduce<
T,
@@ -577,6 +594,7 @@ TYPED_TEST(AsyncReduceTests, TestAsyncReducePolicyInitCustomPlus)
TYPED_TEST(AsyncReduceTests, TestAsyncReducePolicyAllocatorInitCustomPlus)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
testAsyncReduce<
T,
@@ -587,6 +605,7 @@ TYPED_TEST(AsyncReduceTests, TestAsyncReducePolicyAllocatorInitCustomPlus)
TYPED_TEST(AsyncReduceTests, TestAsyncReducePolicyOnInitCustomPlus)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
testAsyncReduce<
T,
@@ -597,6 +616,7 @@ TYPED_TEST(AsyncReduceTests, TestAsyncReducePolicyOnInitCustomPlus)
TYPED_TEST(AsyncReduceTests, TestAsyncReducePolicyAllocatorOnInitCustomPlus)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
testAsyncReduce<
T,
@@ -658,6 +678,7 @@ void testAsyncReduceCountingIterator()
TYPED_TEST(AsyncReduceTests, TestAsyncReduceCountingIterator)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
testAsyncReduceCountingIterator<
T,
@@ -668,6 +689,7 @@ TYPED_TEST(AsyncReduceTests, TestAsyncReduceCountingIterator)
TYPED_TEST(AsyncReduceTests, TestAsyncReducePolicyCountingIterator)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
testAsyncReduceCountingIterator<
T,
@@ -678,6 +700,7 @@ TYPED_TEST(AsyncReduceTests, TestAsyncReducePolicyCountingIterator)
TYPED_TEST(AsyncReduceTests, TestAsyncReduceCountingIteratorInit)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
testAsyncReduceCountingIterator<
T,
@@ -688,6 +711,7 @@ TYPED_TEST(AsyncReduceTests, TestAsyncReduceCountingIteratorInit)
TYPED_TEST(AsyncReduceTests, TestAsyncReducePolicyCountingIteratorInit)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
testAsyncReduceCountingIterator<
T,
@@ -698,6 +722,7 @@ TYPED_TEST(AsyncReduceTests, TestAsyncReducePolicyCountingIteratorInit)
TYPED_TEST(AsyncReduceTests, TestAsyncReduceCountingIteratorInitPlus)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
testAsyncReduceCountingIterator<
T,
@@ -708,6 +733,7 @@ TYPED_TEST(AsyncReduceTests, TestAsyncReduceCountingIteratorInitPlus)
TYPED_TEST(AsyncReduceTests, TestAsyncReducePolicyCountingIteratorInitPlus)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
testAsyncReduceCountingIterator<
T,
@@ -718,6 +744,7 @@ TYPED_TEST(AsyncReduceTests, TestAsyncReducePolicyCountingIteratorInitPlus)
TYPED_TEST(AsyncReduceTests, TestAsyncReduceCountingIteratorInitCustomPlus)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
testAsyncReduceCountingIterator<
T,
@@ -728,6 +755,7 @@ TYPED_TEST(AsyncReduceTests, TestAsyncReduceCountingIteratorInitCustomPlus)
TYPED_TEST(AsyncReduceTests, TestAsyncReducePolicyCountingIteratorInitCustomPlus)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
testAsyncReduceCountingIterator<
T,
@@ -740,6 +768,7 @@ TYPED_TEST(AsyncReduceTests, TestAsyncReducePolicyCountingIteratorInitCustomPlus
TYPED_TEST(AsyncReduceTests, TestAsyncReduceUsing)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
for(auto size : get_sizes())
{
@@ -796,6 +825,7 @@ TYPED_TEST(AsyncReduceTests, TestAsyncReduceUsing)
TYPED_TEST(AsyncReduceTests, TestAsyncReduceAfter)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
for(auto size : get_sizes())
{
@@ -872,6 +902,7 @@ TYPED_TEST(AsyncReduceTests, TestAsyncReduceAfter)
TYPED_TEST(AsyncReduceTests, TestAsyncReduceOnThenAfter)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
for(auto size : get_sizes())
{
@@ -956,6 +987,7 @@ TYPED_TEST(AsyncReduceTests, TestAsyncReduceOnThenAfter)
/*
TYPED_TEST(AsyncReduceTests, TestAsyncReduceAllocatorOnThenAfter)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
for(auto size : get_sizes())
{
@@ -1042,6 +1074,7 @@ TYPED_TEST(AsyncReduceTests, TestAsyncReduceAllocatorOnThenAfter)
TYPED_TEST(AsyncReduceTests, TestAsyncReduceCaching)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
for(auto size : get_sizes())
{
@@ -1097,6 +1130,7 @@ TYPED_TEST(AsyncReduceTests, TestAsyncReduceCaching)
TYPED_TEST(AsyncReduceTests, TestAsyncCopyThenReduce)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
for(auto size : get_sizes())
{

MODIFY
test/test_async_sort.cpp
test/test_async_sort.cpp
@@ -227,78 +227,91 @@ void TestAsyncSort()
TYPED_TEST(AsyncSortTests, AsyncSort)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
TestAsyncSort<T, sort_invoker, wait_for_futures>();
};
TYPED_TEST(AsyncSortTests, AsyncSortNoWait)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
TestAsyncSort<T, sort_invoker, do_not_wait_for_futures>();
};
TYPED_TEST(AsyncSortTests, AsyncSortPolicy)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
TestAsyncSort<T, sort_invoker_device, wait_for_futures>();
};
TYPED_TEST(AsyncSortTests, AsyncSortLess)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
TestAsyncSort<T, sort_invoker_less, wait_for_futures>();
};
TYPED_TEST(AsyncSortTests, AsyncSortLessNoWait)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
TestAsyncSort<T, sort_invoker_less, do_not_wait_for_futures>();
};
TYPED_TEST(AsyncSortTests, AsyncSortPloicyLess)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
TestAsyncSort<T, sort_invoker_less_device, wait_for_futures>();
};
TYPED_TEST(AsyncSortTests, AsyncSortGreater)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
TestAsyncSort<T, sort_invoker_greater, wait_for_futures>();
};
TYPED_TEST(AsyncSortTests, AsyncSortGreaterNoWait)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
TestAsyncSort<T, sort_invoker_greater, do_not_wait_for_futures>();
};
TYPED_TEST(AsyncSortTests, AsyncSortPolicyGreaterNoWait)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
TestAsyncSort<T, sort_invoker_greater_device, do_not_wait_for_futures>();
};
TYPED_TEST(AsyncSortTests, AsyncSortCustomGreater)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
TestAsyncSort<T, sort_invoker_custom_greater, wait_for_futures>();
};
TYPED_TEST(AsyncSortTests, AsyncSortCustomGreaterNoWait)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
TestAsyncSort<T, sort_invoker_custom_greater, do_not_wait_for_futures>();
};
TYPED_TEST(AsyncSortTests, AsyncSortPolicyCustomGreater)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
TestAsyncSort<T, sort_invoker_custom_greater_device, wait_for_futures>();
};
TYPED_TEST(AsyncSortTests, AsyncSortPolicyCustomGreaterNoWait)
{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
using T = typename TestFixture::input_type;
TestAsyncSort<T, sort_invoker_custom_greater_device, do_not_wait_for_futures>();
};

ADD
test/test_async_transform.cpp
+#include <thrust/detail/config.h>
+
+#if THRUST_CPP_DIALECT >= 2011 && !defined(THRUST_LEGACY_GCC)
+
+#include <thrust/async/transform.h>
+#include <thrust/async/copy.h>
+#include <thrust/host_vector.h>
+#include <thrust/device_vector.h>
+
+#include "test_header.hpp"
+
+template <typename T>
+struct divide_by_2
+{
+ __host__ __device__
+ T operator()(T x) const
+ {
+ return x / 2;
+ }
+};
+
+#define DEFINE_STATEFUL_ASYNC_TRANSFORM_UNARY_INVOKER( \
+ NAME, MEMBERS, CTOR, DTOR, VALIDATE, ... \
+ ) \
+ template <typename T> \
+ struct NAME \
+ { \
+ MEMBERS \
+ \
+ NAME() { CTOR } \
+ \
+ ~NAME() { DTOR } \
+ \
+ template <typename Event> \
+ void validate_event(Event& e) \
+ { \
+ THRUST_UNUSED_VAR(e); \
+ VALIDATE \
+ } \
+ \
+ template < \
+ typename ForwardIt, typename Sentinel, typename OutputIt \
+ , typename UnaryOperation \
+ > \
+ __host__ \
+ auto operator()( \
+ ForwardIt&& first, Sentinel&& last, OutputIt&& output \
+ , UnaryOperation&& op \
+ ) \
+ THRUST_DECLTYPE_RETURNS( \
+ ::thrust::async::transform( \
+ __VA_ARGS__ \
+ ) \
+ ) \
+ }; \
+ /**/
+
+#define DEFINE_ASYNC_TRANSFORM_UNARY_INVOKER(NAME, ...) \
+ DEFINE_STATEFUL_ASYNC_TRANSFORM_UNARY_INVOKER( \
+ NAME \
+ , THRUST_PP_EMPTY(), THRUST_PP_EMPTY(), THRUST_PP_EMPTY(), THRUST_PP_EMPTY()\
+ , __VA_ARGS__ \
+ ) \
+ /**/
+
+#define DEFINE_SYNC_TRANSFORM_UNARY_INVOKER(NAME, ...) \
+ template <typename T> \
+ struct NAME \
+ { \
+ \
+ template < \
+ typename ForwardIt, typename Sentinel, typename OutputIt \
+ , typename UnaryOperation \
+ > \
+ __host__ \
+ auto operator()( \
+ ForwardIt&& first, Sentinel&& last, OutputIt&& output \
+ , UnaryOperation&& op \
+ ) \
+ THRUST_DECLTYPE_RETURNS( \
+ ::thrust::transform( \
+ __VA_ARGS__ \
+ ) \
+ ) \
+ }; \
+ /**/
+
+DEFINE_ASYNC_TRANSFORM_UNARY_INVOKER(
+ transform_unary_async_invoker
+, THRUST_FWD(first), THRUST_FWD(last)
+, THRUST_FWD(output)
+, THRUST_FWD(op)
+);
+DEFINE_ASYNC_TRANSFORM_UNARY_INVOKER(
+ transform_unary_async_invoker_device
+, thrust::device
+, THRUST_FWD(first), THRUST_FWD(last)
+, THRUST_FWD(output)
+, THRUST_FWD(op)
+);
+DEFINE_ASYNC_TRANSFORM_UNARY_INVOKER(
+ transform_unary_async_invoker_device_allocator
+, thrust::device(thrust::device_allocator<void>{})
+, THRUST_FWD(first), THRUST_FWD(last)
+, THRUST_FWD(output)
+, THRUST_FWD(op)
+);
+DEFINE_STATEFUL_ASYNC_TRANSFORM_UNARY_INVOKER(
+ transform_unary_async_invoker_device_on
+ // Members.
+, hipStream_t stream_;
+ // Constructor.
+, thrust::hip_rocprim::throw_on_error(
+ hipStreamCreateWithFlags(&stream_, hipStreamNonBlocking)
+ );
+ // Destructor.
+, thrust::hip_rocprim::throw_on_error(
+ hipStreamDestroy(stream_)
+ );
+ // `validate_event` member.
+, ASSERT_EQ_QUIET(stream_, e.stream().native_handle());
+ // Arguments to `thrust::async::transform`.
+, thrust::device.on(stream_)
+, THRUST_FWD(first), THRUST_FWD(last)
+, THRUST_FWD(output)
+, THRUST_FWD(op)
+);
+DEFINE_STATEFUL_ASYNC_TRANSFORM_UNARY_INVOKER(
+ transform_unary_async_invoker_device_allocator_on
+ // Members.
+, hipStream_t stream_;
+ // Constructor.
+, thrust::hip_rocprim::throw_on_error(
+ hipStreamCreateWithFlags(&stream_, hipStreamNonBlocking)
+ );
+ // Destructor.
+, thrust::hip_rocprim::throw_on_error(
+ hipStreamDestroy(stream_)
+ );
+ // `validate_event` member.
+, ASSERT_EQ_QUIET(stream_, e.stream().native_handle());
+ // Arguments to `thrust::async::transform`.
+, thrust::device(thrust::device_allocator<void>{}).on(stream_)
+, THRUST_FWD(first), THRUST_FWD(last)
+, THRUST_FWD(output)
+, THRUST_FWD(op)
+);
+
+DEFINE_SYNC_TRANSFORM_UNARY_INVOKER(
+ transform_unary_sync_invoker
+, THRUST_FWD(first), THRUST_FWD(last)
+, THRUST_FWD(output)
+, THRUST_FWD(op)
+);
+
+///////////////////////////////////////////////////////////////////////////////
+
+TESTS_DEFINE(AsyncTransformTests, NumericalTestsParams);
+
+template <
+ typename T
+, template <typename> class AsyncTransformUnaryInvoker
+, template <typename> class SyncTransformUnaryInvoker
+, template <typename> class UnaryOperation
+>
+void test_async_transform_unary()
+{
+ for(auto size : get_sizes())
+ {
+ SCOPED_TRACE(testing::Message() << "with size = " << size);
+ for(size_t seed_index = 0; seed_index < random_seeds_count + seed_size; seed_index++)
+ {
+ unsigned int seed_value
+ = seed_index < random_seeds_count ? rand() : seeds[seed_index - random_seeds_count];
+ SCOPED_TRACE(testing::Message() << "with seed= " << seed_value);
+
+ thrust::host_vector<T> h0 = get_random_data<T>(size, T(-1000), T(1000), seed_value);
+
+ thrust::device_vector<T> d0a(h0);
+ thrust::device_vector<T> d0b(h0);
+ thrust::device_vector<T> d0c(h0);
+ thrust::device_vector<T> d0d(h0);
+
+ thrust::host_vector<T> h1(size);
+
+ thrust::device_vector<T> d1a(size);
+ thrust::device_vector<T> d1b(size);
+ thrust::device_vector<T> d1c(size);
+ thrust::device_vector<T> d1d(size);
+
+ AsyncTransformUnaryInvoker<T> invoke_async;
+ SyncTransformUnaryInvoker<T> invoke_sync;
+
+ UnaryOperation<T> op;
+
+ ASSERT_EQ(h0, d0a);
+ ASSERT_EQ(h0, d0b);
+ ASSERT_EQ(h0, d0c);
+ ASSERT_EQ(h0, d0d);
+
+ auto f0a = invoke_async(d0a.begin(), d0a.end(), d1a.begin(), op);
+ auto f0b = invoke_async(d0b.begin(), d0b.end(), d1b.begin(), op);
+ auto f0c = invoke_async(d0c.begin(), d0c.end(), d1c.begin(), op);
+ auto f0d = invoke_async(d0d.begin(), d0d.end(), d1d.begin(), op);
+
+ invoke_async.validate_event(f0a);
+ invoke_async.validate_event(f0b);
+ invoke_async.validate_event(f0c);
+ invoke_async.validate_event(f0d);
+
+ // This potentially runs concurrently with the copies.
+ invoke_sync(h0.begin(), h0.end(), h1.begin(), op);
+
+ //TEST_EVENT_WAIT(thrust::when_all(f0a, f0b, f0c, f0d));
+
+ ASSERT_EQ(h0, d0a);
+ ASSERT_EQ(h0, d0b);
+ ASSERT_EQ(h0, d0c);
+ ASSERT_EQ(h0, d0d);
+
+ ASSERT_EQ(h1, d1a);
+ ASSERT_EQ(h1, d1b);
+ ASSERT_EQ(h1, d1c);
+ ASSERT_EQ(h1, d1d);
+ }
+ }
+};
+
+TYPED_TEST(AsyncTransformTests, TestAsyncTransformUnary)
+{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
+ using T = typename TestFixture::input_type;
+ test_async_transform_unary<
+ T,
+ transform_unary_async_invoker
+ , transform_unary_sync_invoker
+ , divide_by_2
+ >();
+}
+
+TYPED_TEST(AsyncTransformTests, TestAsyncTransformUnaryDevice)
+{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
+ using T = typename TestFixture::input_type;
+ test_async_transform_unary<
+ T,
+ transform_unary_async_invoker_device
+ , transform_unary_sync_invoker
+ , divide_by_2
+ >();
+}
+
+TYPED_TEST(AsyncTransformTests, TestAsyncTransformUnaryDeviceAllocator)
+{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
+ using T = typename TestFixture::input_type;
+ test_async_transform_unary<
+ T,
+ transform_unary_async_invoker_device_allocator
+ , transform_unary_sync_invoker
+ , divide_by_2
+ >();
+}
+
+TYPED_TEST(AsyncTransformTests, TestAsyncTransformUnaryDeviceOn)
+{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
+ using T = typename TestFixture::input_type;
+ test_async_transform_unary<
+ T,
+ transform_unary_async_invoker_device_on
+ , transform_unary_sync_invoker
+ , divide_by_2
+ >();
+}
+
+TYPED_TEST(AsyncTransformTests, TestAsyncTransformUnaryDeviceAllocatorOn)
+{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
+ using T = typename TestFixture::input_type;
+ test_async_transform_unary<
+ T,
+ transform_unary_async_invoker_device_allocator_on
+ , transform_unary_sync_invoker
+ , divide_by_2
+ >();
+}
+
+///////////////////////////////////////////////////////////////////////////////
+
+template <
+ typename T
+, template <typename> class AsyncTransformUnaryInvoker
+, template <typename> class SyncTransformUnaryInvoker
+, template <typename> class UnaryOperation
+>
+void test_async_transform_unary_inplace()
+{
+ for(auto size : get_sizes())
+ {
+ SCOPED_TRACE(testing::Message() << "with size = " << size);
+ for(size_t seed_index = 0; seed_index < random_seeds_count + seed_size; seed_index++)
+ {
+ unsigned int seed_value
+ = seed_index < random_seeds_count ? rand() : seeds[seed_index - random_seeds_count];
+ SCOPED_TRACE(testing::Message() << "with seed= " << seed_value);
+
+ thrust::host_vector<T> h0 = get_random_data<T>(size, T(-1000), T(1000), seed_value);
+
+ thrust::device_vector<T> d0a(h0);
+ thrust::device_vector<T> d0b(h0);
+ thrust::device_vector<T> d0c(h0);
+ thrust::device_vector<T> d0d(h0);
+
+ AsyncTransformUnaryInvoker<T> invoke_async;
+ SyncTransformUnaryInvoker<T> invoke_sync;
+
+ UnaryOperation<T> op;
+
+ ASSERT_EQ(h0, d0a);
+ ASSERT_EQ(h0, d0b);
+ ASSERT_EQ(h0, d0c);
+ ASSERT_EQ(h0, d0d);
+
+ auto f0a = invoke_async(d0a.begin(), d0a.end(), d0a.begin(), op);
+ auto f0b = invoke_async(d0b.begin(), d0b.end(), d0b.begin(), op);
+ auto f0c = invoke_async(d0c.begin(), d0c.end(), d0c.begin(), op);
+ auto f0d = invoke_async(d0d.begin(), d0d.end(), d0d.begin(), op);
+
+ invoke_async.validate_event(f0a);
+ invoke_async.validate_event(f0b);
+ invoke_async.validate_event(f0c);
+ invoke_async.validate_event(f0d);
+
+ // This potentially runs concurrently with the copies.
+ invoke_sync(h0.begin(), h0.end(), h0.begin(), op);
+
+ //TEST_EVENT_WAIT(thrust::when_all(f0a, f0b, f0c, f0d));
+
+ ASSERT_EQ(h0, d0a);
+ ASSERT_EQ(h0, d0b);
+ ASSERT_EQ(h0, d0c);
+ ASSERT_EQ(h0, d0d);
+ }
+ }
+};
+
+TYPED_TEST(AsyncTransformTests, TestAsyncTransformUnaryInplace)
+{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
+ using T = typename TestFixture::input_type;
+ test_async_transform_unary_inplace<
+ T,
+ transform_unary_async_invoker
+ , transform_unary_sync_invoker
+ , divide_by_2
+ >();
+}
+
+TYPED_TEST(AsyncTransformTests, TestAsyncTransformUnaryInplaceDevice)
+{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
+ using T = typename TestFixture::input_type;
+ test_async_transform_unary_inplace<
+ T,
+ transform_unary_async_invoker_device
+ , transform_unary_sync_invoker
+ , divide_by_2
+ >();
+}
+
+TYPED_TEST(AsyncTransformTests, TestAsyncTransformUnaryInplaceDeviceAllocator)
+{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
+ using T = typename TestFixture::input_type;
+ test_async_transform_unary_inplace<
+ T,
+ transform_unary_async_invoker_device_allocator
+ , transform_unary_sync_invoker
+ , divide_by_2
+ >();
+}
+
+TYPED_TEST(AsyncTransformTests, TestAsyncTransformUnaryInplaceDeviceOn)
+{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
+ using T = typename TestFixture::input_type;
+ test_async_transform_unary_inplace<
+ T,
+ transform_unary_async_invoker_device_on
+ , transform_unary_sync_invoker
+ , divide_by_2
+ >();
+}
+
+TYPED_TEST(AsyncTransformTests, TestAsyncTransformUnaryInplaceDeviceAllocatorOn)
+{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
+ using T = typename TestFixture::input_type;
+ test_async_transform_unary_inplace<
+ T,
+ transform_unary_async_invoker_device_allocator_on
+ , transform_unary_sync_invoker
+ , divide_by_2
+ >();
+}
+
+///////////////////////////////////////////////////////////////////////////////
+
+#endif

MODIFY
thrust/async/transform.h
thrust/async/transform.h
@@ -53,6 +53,7 @@ async_transform(
, ForwardIt first, Sentinel last, OutputIt output, UnaryOperation op
)
{
+ (void) exec; (void) first; (void) last; (void) output; (void) op;
THRUST_STATIC_ASSERT_MSG(
(thrust::detail::depend_on_instantiation<ForwardIt, false>::value)
, "this algorithm is not implemented for the specified system"

Added transform async tests.