Centralize the fix for testing folder, fix async reduce into and allocator policies
MODIFY
testing/allocator_aware_policies.cu
testing/allocator_aware_policies.cu
#include <unittest/unittest.h>
#include <thrust/detail/seq.h>
-#include <thrust/system/cpp/detail/par.h>
+
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HIP
+ #include <thrust/system/hip/detail/par.h>
+#elif THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_NVCC
#include <thrust/system/cuda/detail/par.h>
+#endif
+
+#include <thrust/system/cpp/detail/par.h>
#include <thrust/system/omp/detail/par.h>
#include <thrust/system/tbb/detail/par.h>
@@ -121,9 +127,9 @@ typedef policy_info<
thrust::system::cpp::detail::execution_policy
> cpp_par_info;
typedef policy_info<
- thrust::system::cuda::detail::par_t,
- thrust::cuda_cub::execute_on_stream_base
-> cuda_par_info;
+ thrust::system::THRUST_DEVICE_BACKEND::detail::par_t,
+ thrust::THRUST_DEVICE_BACKEND_DETAIL::execute_on_stream_base
+> THRUST_DEVICE_BACKEND_par_info;
typedef policy_info<
thrust::system::omp::detail::par_t,
thrust::system::omp::detail::execution_policy
@@ -138,7 +144,7 @@ SimpleUnitTest<
unittest::type_list<
sequential_info,
cpp_par_info,
- cuda_par_info,
+ THRUST_DEVICE_BACKEND_par_info,
omp_par_info,
tbb_par_info
>

MODIFY
testing/async_reduce.cu
testing/async_reduce.cu
@@ -101,14 +101,14 @@ DEFINE_ASYNC_REDUCE_INVOKER(
DEFINE_STATEFUL_ASYNC_REDUCE_INVOKER(
reduce_async_invoker_device_on
// Members.
-, cudaStream_t stream_;
+, SPECIALIZE_DEVICE_RESOURCE_NAME(Stream_t) stream_;
// Constructor.
-, thrust::cuda_cub::throw_on_error(
- cudaStreamCreateWithFlags(&stream_, cudaStreamNonBlocking)
+, thrust::THRUST_DEVICE_BACKEND_DETAIL::throw_on_error(
+ SPECIALIZE_DEVICE_RESOURCE_NAME(StreamCreateWithFlags)(&stream_, SPECIALIZE_DEVICE_RESOURCE_NAME(StreamNonBlocking))
);
// Destructor.
-, thrust::cuda_cub::throw_on_error(
- cudaStreamDestroy(stream_)
+, thrust::THRUST_DEVICE_BACKEND_DETAIL::throw_on_error(
+ SPECIALIZE_DEVICE_RESOURCE_NAME(StreamDestroy)(stream_)
);
// `validate_event` member.
, ASSERT_EQUAL_QUIET(stream_, e.stream().native_handle());
@@ -119,14 +119,14 @@ DEFINE_STATEFUL_ASYNC_REDUCE_INVOKER(
DEFINE_STATEFUL_ASYNC_REDUCE_INVOKER(
reduce_async_invoker_device_allocator_on
// Members.
-, cudaStream_t stream_;
+, SPECIALIZE_DEVICE_RESOURCE_NAME(Stream_t) stream_;
// Constructor.
-, thrust::cuda_cub::throw_on_error(
- cudaStreamCreateWithFlags(&stream_, cudaStreamNonBlocking)
+, thrust::THRUST_DEVICE_BACKEND_DETAIL::throw_on_error(
+ SPECIALIZE_DEVICE_RESOURCE_NAME(StreamCreateWithFlags)(&stream_, SPECIALIZE_DEVICE_RESOURCE_NAME(StreamNonBlocking))
);
// Destructor.
-, thrust::cuda_cub::throw_on_error(
- cudaStreamDestroy(stream_)
+, thrust::THRUST_DEVICE_BACKEND_DETAIL::throw_on_error(
+ SPECIALIZE_DEVICE_RESOURCE_NAME(StreamDestroy)(stream_)
);
// `validate_event` member.
, ASSERT_EQUAL_QUIET(stream_, e.stream().native_handle());
@@ -160,14 +160,14 @@ DEFINE_ASYNC_REDUCE_INVOKER(
DEFINE_STATEFUL_ASYNC_REDUCE_INVOKER(
reduce_async_invoker_device_on_init
// Members.
-, cudaStream_t stream_;
+, SPECIALIZE_DEVICE_RESOURCE_NAME(Stream_t) stream_;
// Constructor.
-, thrust::cuda_cub::throw_on_error(
- cudaStreamCreateWithFlags(&stream_, cudaStreamNonBlocking)
+, thrust::THRUST_DEVICE_BACKEND_DETAIL::throw_on_error(
+ SPECIALIZE_DEVICE_RESOURCE_NAME(StreamCreateWithFlags)(&stream_, SPECIALIZE_DEVICE_RESOURCE_NAME(StreamNonBlocking))
);
// Destructor.
-, thrust::cuda_cub::throw_on_error(
- cudaStreamDestroy(stream_)
+, thrust::THRUST_DEVICE_BACKEND_DETAIL::throw_on_error(
+ SPECIALIZE_DEVICE_RESOURCE_NAME(StreamDestroy)(stream_)
);
// `validate_event` member.
, ASSERT_EQUAL_QUIET(stream_, e.stream().native_handle());
@@ -179,14 +179,14 @@ DEFINE_STATEFUL_ASYNC_REDUCE_INVOKER(
DEFINE_STATEFUL_ASYNC_REDUCE_INVOKER(
reduce_async_invoker_device_allocator_on_init
// Members.
-, cudaStream_t stream_;
+, SPECIALIZE_DEVICE_RESOURCE_NAME(Stream_t) stream_;
// Constructor.
-, thrust::cuda_cub::throw_on_error(
- cudaStreamCreateWithFlags(&stream_, cudaStreamNonBlocking)
+, thrust::THRUST_DEVICE_BACKEND_DETAIL::throw_on_error(
+ SPECIALIZE_DEVICE_RESOURCE_NAME(StreamCreateWithFlags)(&stream_, SPECIALIZE_DEVICE_RESOURCE_NAME(StreamNonBlocking))
);
// Destructor.
-, thrust::cuda_cub::throw_on_error(
- cudaStreamDestroy(stream_)
+, thrust::THRUST_DEVICE_BACKEND_DETAIL::throw_on_error(
+ SPECIALIZE_DEVICE_RESOURCE_NAME(StreamDestroy)(stream_)
);
// `validate_event` member.
, ASSERT_EQUAL_QUIET(stream_, e.stream().native_handle());
@@ -225,14 +225,14 @@ DEFINE_ASYNC_REDUCE_INVOKER(
DEFINE_STATEFUL_ASYNC_REDUCE_INVOKER(
reduce_async_invoker_device_on_init_plus
// Members.
-, cudaStream_t stream_;
+, SPECIALIZE_DEVICE_RESOURCE_NAME(Stream_t) stream_;
// Constructor.
-, thrust::cuda_cub::throw_on_error(
- cudaStreamCreateWithFlags(&stream_, cudaStreamNonBlocking)
+, thrust::THRUST_DEVICE_BACKEND_DETAIL::throw_on_error(
+ SPECIALIZE_DEVICE_RESOURCE_NAME(StreamCreateWithFlags)(&stream_, SPECIALIZE_DEVICE_RESOURCE_NAME(StreamNonBlocking))
);
// Destructor.
-, thrust::cuda_cub::throw_on_error(
- cudaStreamDestroy(stream_)
+, thrust::THRUST_DEVICE_BACKEND_DETAIL::throw_on_error(
+ SPECIALIZE_DEVICE_RESOURCE_NAME(StreamDestroy)(stream_)
);
// `validate_event` member.
, ASSERT_EQUAL_QUIET(stream_, e.stream().native_handle());
@@ -245,14 +245,14 @@ DEFINE_STATEFUL_ASYNC_REDUCE_INVOKER(
DEFINE_STATEFUL_ASYNC_REDUCE_INVOKER(
reduce_async_invoker_device_allocator_on_init_plus
// Members.
-, cudaStream_t stream_;
+, SPECIALIZE_DEVICE_RESOURCE_NAME(Stream_t) stream_;
// Constructor.
-, thrust::cuda_cub::throw_on_error(
- cudaStreamCreateWithFlags(&stream_, cudaStreamNonBlocking)
+, thrust::THRUST_DEVICE_BACKEND_DETAIL::throw_on_error(
+ SPECIALIZE_DEVICE_RESOURCE_NAME(StreamCreateWithFlags)(&stream_, SPECIALIZE_DEVICE_RESOURCE_NAME(StreamNonBlocking))
);
// Destructor.
-, thrust::cuda_cub::throw_on_error(
- cudaStreamDestroy(stream_)
+, thrust::THRUST_DEVICE_BACKEND_DETAIL::throw_on_error(
+ SPECIALIZE_DEVICE_RESOURCE_NAME(StreamDestroy)(stream_)
);
// `validate_event` member.
, ASSERT_EQUAL_QUIET(stream_, e.stream().native_handle());
@@ -293,14 +293,14 @@ DEFINE_ASYNC_REDUCE_INVOKER(
DEFINE_STATEFUL_ASYNC_REDUCE_INVOKER(
reduce_async_invoker_device_on_init_custom_plus
// Members.
-, cudaStream_t stream_;
+, SPECIALIZE_DEVICE_RESOURCE_NAME(Stream_t) stream_;
// Constructor.
-, thrust::cuda_cub::throw_on_error(
- cudaStreamCreateWithFlags(&stream_, cudaStreamNonBlocking)
+, thrust::THRUST_DEVICE_BACKEND_DETAIL::throw_on_error(
+ SPECIALIZE_DEVICE_RESOURCE_NAME(StreamCreateWithFlags)(&stream_, SPECIALIZE_DEVICE_RESOURCE_NAME(StreamNonBlocking))
);
// Destructor.
-, thrust::cuda_cub::throw_on_error(
- cudaStreamDestroy(stream_)
+, thrust::THRUST_DEVICE_BACKEND_DETAIL::throw_on_error(
+ SPECIALIZE_DEVICE_RESOURCE_NAME(StreamDestroy)(stream_)
);
// `validate_event` member.
, ASSERT_EQUAL_QUIET(stream_, e.stream().native_handle());
@@ -313,14 +313,14 @@ DEFINE_STATEFUL_ASYNC_REDUCE_INVOKER(
DEFINE_STATEFUL_ASYNC_REDUCE_INVOKER(
reduce_async_invoker_device_allocator_on_init_custom_plus
// Members.
-, cudaStream_t stream_;
+, SPECIALIZE_DEVICE_RESOURCE_NAME(Stream_t) stream_;
// Constructor.
-, thrust::cuda_cub::throw_on_error(
- cudaStreamCreateWithFlags(&stream_, cudaStreamNonBlocking)
+, thrust::THRUST_DEVICE_BACKEND_DETAIL::throw_on_error(
+ SPECIALIZE_DEVICE_RESOURCE_NAME(StreamCreateWithFlags)(&stream_, SPECIALIZE_DEVICE_RESOURCE_NAME(StreamNonBlocking))
);
// Destructor.
-, thrust::cuda_cub::throw_on_error(
- cudaStreamDestroy(stream_)
+, thrust::THRUST_DEVICE_BACKEND_DETAIL::throw_on_error(
+ SPECIALIZE_DEVICE_RESOURCE_NAME(StreamDestroy)(stream_)
);
// `validate_event` member.
, ASSERT_EQUAL_QUIET(stream_, e.stream().native_handle());
@@ -848,9 +848,9 @@ struct test_async_reduce_on_then_after
ASSERT_EQUAL(h0, d0);
- cudaStream_t stream;
- thrust::cuda_cub::throw_on_error(
- cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking)
+ SPECIALIZE_DEVICE_RESOURCE_NAME(Stream_t) stream;
+ thrust::THRUST_DEVICE_BACKEND_DETAIL::throw_on_error(
+ SPECIALIZE_DEVICE_RESOURCE_NAME(StreamCreateWithFlags)(&stream, SPECIALIZE_DEVICE_RESOURCE_NAME(StreamNonBlocking))
);
auto f0 = thrust::async::reduce(
@@ -900,8 +900,8 @@ struct test_async_reduce_on_then_after
ASSERT_EQUAL(r0, r1);
- thrust::cuda_cub::throw_on_error(
- cudaStreamDestroy(stream)
+ thrust::THRUST_DEVICE_BACKEND_DETAIL::throw_on_error(
+ SPECIALIZE_DEVICE_RESOURCE_NAME(StreamDestroy)(stream)
);
}
};
@@ -923,14 +923,14 @@ struct test_async_reduce_allocator_on_then_after
ASSERT_EQUAL(h0, d0);
- cudaStream_t stream0;
- thrust::cuda_cub::throw_on_error(
- cudaStreamCreateWithFlags(&stream0, cudaStreamNonBlocking)
+ SPECIALIZE_DEVICE_RESOURCE_NAME(Stream_t) stream0;
+ thrust::THRUST_DEVICE_BACKEND_DETAIL::throw_on_error(
+ SPECIALIZE_DEVICE_RESOURCE_NAME(StreamCreateWithFlags)(&stream0, SPECIALIZE_DEVICE_RESOURCE_NAME(StreamNonBlocking))
);
- cudaStream_t stream1;
- thrust::cuda_cub::throw_on_error(
- cudaStreamCreateWithFlags(&stream1, cudaStreamNonBlocking)
+ SPECIALIZE_DEVICE_RESOURCE_NAME(Stream_t) stream1;
+ thrust::THRUST_DEVICE_BACKEND_DETAIL::throw_on_error(
+ SPECIALIZE_DEVICE_RESOURCE_NAME(StreamCreateWithFlags)(&stream1, SPECIALIZE_DEVICE_RESOURCE_NAME(StreamNonBlocking))
);
auto f0 = thrust::async::reduce(
@@ -984,8 +984,8 @@ struct test_async_reduce_allocator_on_then_after
ASSERT_EQUAL(r0, r1);
- thrust::cuda_cub::throw_on_error(cudaStreamDestroy(stream0));
- thrust::cuda_cub::throw_on_error(cudaStreamDestroy(stream1));
+ thrust::THRUST_DEVICE_BACKEND_DETAIL::throw_on_error(SPECIALIZE_DEVICE_RESOURCE_NAME(StreamDestroy)(stream0));
+ thrust::THRUST_DEVICE_BACKEND_DETAIL::throw_on_error(SPECIALIZE_DEVICE_RESOURCE_NAME(StreamDestroy)(stream1));
}
};
DECLARE_GENERIC_SIZED_UNITTEST_WITH_TYPES(

MODIFY
testing/async_reduce_into.cu
testing/async_reduce_into.cu
@@ -105,14 +105,14 @@ DEFINE_ASYNC_REDUCE_INTO_INVOKER(
DEFINE_STATEFUL_ASYNC_REDUCE_INTO_INVOKER(
reduce_into_async_invoker_device_on
// Members.
-, cudaStream_t stream_;
+, SPECIALIZE_DEVICE_RESOURCE_NAME(Stream_t) stream_;
// Constructor.
-, thrust::cuda_cub::throw_on_error(
- cudaStreamCreateWithFlags(&stream_, cudaStreamNonBlocking)
+, thrust::THRUST_DEVICE_BACKEND_DETAIL::throw_on_error(
+ SPECIALIZE_DEVICE_RESOURCE_NAME(StreamCreateWithFlags)(&stream_, SPECIALIZE_DEVICE_RESOURCE_NAME(StreamNonBlocking))
);
// Destructor.
-, thrust::cuda_cub::throw_on_error(
- cudaStreamDestroy(stream_)
+, thrust::THRUST_DEVICE_BACKEND_DETAIL::throw_on_error(
+ SPECIALIZE_DEVICE_RESOURCE_NAME(StreamDestroy)(stream_)
);
// `validate_event` member.
, ASSERT_EQUAL_QUIET(stream_, e.stream().native_handle());
@@ -124,14 +124,14 @@ DEFINE_STATEFUL_ASYNC_REDUCE_INTO_INVOKER(
DEFINE_STATEFUL_ASYNC_REDUCE_INTO_INVOKER(
reduce_into_async_invoker_device_allocator_on
// Members.
-, cudaStream_t stream_;
+, SPECIALIZE_DEVICE_RESOURCE_NAME(Stream_t) stream_;
// Constructor.
-, thrust::cuda_cub::throw_on_error(
- cudaStreamCreateWithFlags(&stream_, cudaStreamNonBlocking)
+, thrust::THRUST_DEVICE_BACKEND_DETAIL::throw_on_error(
+ SPECIALIZE_DEVICE_RESOURCE_NAME(StreamCreateWithFlags)(&stream_, SPECIALIZE_DEVICE_RESOURCE_NAME(StreamNonBlocking))
);
// Destructor.
-, thrust::cuda_cub::throw_on_error(
- cudaStreamDestroy(stream_)
+, thrust::THRUST_DEVICE_BACKEND_DETAIL::throw_on_error(
+ SPECIALIZE_DEVICE_RESOURCE_NAME(StreamDestroy)(stream_)
);
// `validate_event` member.
, ASSERT_EQUAL_QUIET(stream_, e.stream().native_handle());
@@ -169,14 +169,14 @@ DEFINE_ASYNC_REDUCE_INTO_INVOKER(
DEFINE_STATEFUL_ASYNC_REDUCE_INTO_INVOKER(
reduce_into_async_invoker_device_on_init
// Members.
-, cudaStream_t stream_;
+, SPECIALIZE_DEVICE_RESOURCE_NAME(Stream_t) stream_;
// Constructor.
-, thrust::cuda_cub::throw_on_error(
- cudaStreamCreateWithFlags(&stream_, cudaStreamNonBlocking)
+, thrust::THRUST_DEVICE_BACKEND_DETAIL::throw_on_error(
+ SPECIALIZE_DEVICE_RESOURCE_NAME(StreamCreateWithFlags)(&stream_, SPECIALIZE_DEVICE_RESOURCE_NAME(StreamNonBlocking))
);
// Destructor.
-, thrust::cuda_cub::throw_on_error(
- cudaStreamDestroy(stream_)
+, thrust::THRUST_DEVICE_BACKEND_DETAIL::throw_on_error(
+ SPECIALIZE_DEVICE_RESOURCE_NAME(StreamDestroy)(stream_)
);
// `validate_event` member.
, ASSERT_EQUAL_QUIET(stream_, e.stream().native_handle());
@@ -189,14 +189,14 @@ DEFINE_STATEFUL_ASYNC_REDUCE_INTO_INVOKER(
DEFINE_STATEFUL_ASYNC_REDUCE_INTO_INVOKER(
reduce_into_async_invoker_device_allocator_on_init
// Members.
-, cudaStream_t stream_;
+, SPECIALIZE_DEVICE_RESOURCE_NAME(Stream_t) stream_;
// Constructor.
-, thrust::cuda_cub::throw_on_error(
- cudaStreamCreateWithFlags(&stream_, cudaStreamNonBlocking)
+, thrust::THRUST_DEVICE_BACKEND_DETAIL::throw_on_error(
+ SPECIALIZE_DEVICE_RESOURCE_NAME(StreamCreateWithFlags)(&stream_, SPECIALIZE_DEVICE_RESOURCE_NAME(StreamNonBlocking))
);
// Destructor.
-, thrust::cuda_cub::throw_on_error(
- cudaStreamDestroy(stream_)
+, thrust::THRUST_DEVICE_BACKEND_DETAIL::throw_on_error(
+ SPECIALIZE_DEVICE_RESOURCE_NAME(StreamDestroy)(stream_)
);
// `validate_event` member.
, ASSERT_EQUAL_QUIET(stream_, e.stream().native_handle());
@@ -239,14 +239,14 @@ DEFINE_ASYNC_REDUCE_INTO_INVOKER(
DEFINE_STATEFUL_ASYNC_REDUCE_INTO_INVOKER(
reduce_into_async_invoker_device_on_init_plus
// Members.
-, cudaStream_t stream_;
+, SPECIALIZE_DEVICE_RESOURCE_NAME(Stream_t) stream_;
// Constructor.
-, thrust::cuda_cub::throw_on_error(
- cudaStreamCreateWithFlags(&stream_, cudaStreamNonBlocking)
+, thrust::THRUST_DEVICE_BACKEND_DETAIL::throw_on_error(
+ SPECIALIZE_DEVICE_RESOURCE_NAME(StreamCreateWithFlags)(&stream_, SPECIALIZE_DEVICE_RESOURCE_NAME(StreamNonBlocking))
);
// Destructor.
-, thrust::cuda_cub::throw_on_error(
- cudaStreamDestroy(stream_)
+, thrust::THRUST_DEVICE_BACKEND_DETAIL::throw_on_error(
+ SPECIALIZE_DEVICE_RESOURCE_NAME(StreamDestroy)(stream_)
);
// `validate_event` member.
, ASSERT_EQUAL_QUIET(stream_, e.stream().native_handle());
@@ -260,14 +260,14 @@ DEFINE_STATEFUL_ASYNC_REDUCE_INTO_INVOKER(
DEFINE_STATEFUL_ASYNC_REDUCE_INTO_INVOKER(
reduce_into_async_invoker_device_allocator_on_init_plus
// Members.
-, cudaStream_t stream_;
+, SPECIALIZE_DEVICE_RESOURCE_NAME(Stream_t) stream_;
// Constructor.
-, thrust::cuda_cub::throw_on_error(
- cudaStreamCreateWithFlags(&stream_, cudaStreamNonBlocking)
+, thrust::THRUST_DEVICE_BACKEND_DETAIL::throw_on_error(
+ SPECIALIZE_DEVICE_RESOURCE_NAME(StreamCreateWithFlags)(&stream_, SPECIALIZE_DEVICE_RESOURCE_NAME(StreamNonBlocking))
);
// Destructor.
-, thrust::cuda_cub::throw_on_error(
- cudaStreamDestroy(stream_)
+, thrust::THRUST_DEVICE_BACKEND_DETAIL::throw_on_error(
+ SPECIALIZE_DEVICE_RESOURCE_NAME(StreamDestroy)(stream_)
);
// `validate_event` member.
, ASSERT_EQUAL_QUIET(stream_, e.stream().native_handle());
@@ -312,14 +312,14 @@ DEFINE_ASYNC_REDUCE_INTO_INVOKER(
DEFINE_STATEFUL_ASYNC_REDUCE_INTO_INVOKER(
reduce_into_async_invoker_device_on_init_custom_plus
// Members.
-, cudaStream_t stream_;
+, SPECIALIZE_DEVICE_RESOURCE_NAME(Stream_t) stream_;
// Constructor.
-, thrust::cuda_cub::throw_on_error(
- cudaStreamCreateWithFlags(&stream_, cudaStreamNonBlocking)
+, thrust::THRUST_DEVICE_BACKEND_DETAIL::throw_on_error(
+ SPECIALIZE_DEVICE_RESOURCE_NAME(StreamCreateWithFlags)(&stream_, SPECIALIZE_DEVICE_RESOURCE_NAME(StreamNonBlocking))
);
// Destructor.
-, thrust::cuda_cub::throw_on_error(
- cudaStreamDestroy(stream_)
+, thrust::THRUST_DEVICE_BACKEND_DETAIL::throw_on_error(
+ SPECIALIZE_DEVICE_RESOURCE_NAME(StreamDestroy)(stream_)
);
// `validate_event` member.
, ASSERT_EQUAL_QUIET(stream_, e.stream().native_handle());
@@ -333,14 +333,14 @@ DEFINE_STATEFUL_ASYNC_REDUCE_INTO_INVOKER(
DEFINE_STATEFUL_ASYNC_REDUCE_INTO_INVOKER(
reduce_into_async_invoker_device_allocator_on_init_custom_plus
// Members.
-, cudaStream_t stream_;
+, SPECIALIZE_DEVICE_RESOURCE_NAME(Stream_t) stream_;
// Constructor.
-, thrust::cuda_cub::throw_on_error(
- cudaStreamCreateWithFlags(&stream_, cudaStreamNonBlocking)
+, thrust::THRUST_DEVICE_BACKEND_DETAIL::throw_on_error(
+ SPECIALIZE_DEVICE_RESOURCE_NAME(StreamCreateWithFlags)(&stream_, SPECIALIZE_DEVICE_RESOURCE_NAME(StreamNonBlocking))
);
// Destructor.
-, thrust::cuda_cub::throw_on_error(
- cudaStreamDestroy(stream_)
+, thrust::THRUST_DEVICE_BACKEND_DETAIL::throw_on_error(
+ SPECIALIZE_DEVICE_RESOURCE_NAME(StreamDestroy)(stream_)
);
// `validate_event` member.
, ASSERT_EQUAL_QUIET(stream_, e.stream().native_handle());

MODIFY
testing/dependencies_aware_policies.cu
testing/dependencies_aware_policies.cu
#include <thrust/detail/seq.h>
#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HIP
- #define THRUST_DEVICE_BACKEND hip
- #define THRUST_DEVICE_BACKEND_DETAIL hip_rocprim
#include <thrust/system/hip/detail/par.h>
#elif THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_NVCC
- #define THRUST_DEVICE_BACKEND cuda
- #define THRUST_DEVICE_BACKEND_DETAIL cuda_cub
#include <thrust/system/cuda/detail/par.h>
#endif

MODIFY
testing/unittest/unittest.h
testing/unittest/unittest.h
#include <unittest/testframework.h>
#include <unittest/special_types.h>
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HIP
+ #define THRUST_DEVICE_BACKEND hip
+ #define THRUST_DEVICE_BACKEND_DETAIL hip_rocprim
+ #define SPECIALIZE_DEVICE_RESOURCE_NAME(name) hip##name
+#elif THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_NVCC
+ #define THRUST_DEVICE_BACKEND cuda
+ #define THRUST_DEVICE_BACKEND_DETAIL cuda_cub
+ #define SPECIALIZE_DEVICE_RESOURCE_NAME(name) cuda##name
+#endif
+

Centralize the fix for testing folder, fix async reduce into and allocator policies