Fixing async reduce for hip and enabling test.
MODIFY
test/test_async_reduce.cpp
test/test_async_reduce.cpp
@@ -142,19 +142,19 @@ DEFINE_SYNC_REDUCE_INVOKER(
DEFINE_ASYNC_REDUCE_INVOKER(
reduce_async_invoker_init
, THRUST_FWD(first), THRUST_FWD(last)
-, unittest::random_integer<T>()
+, random_integer<T>()
);
DEFINE_ASYNC_REDUCE_INVOKER(
reduce_async_invoker_device_init
, thrust::device
, THRUST_FWD(first), THRUST_FWD(last)
-, unittest::random_integer<T>()
+, random_integer<T>()
);
DEFINE_ASYNC_REDUCE_INVOKER(
reduce_async_invoker_device_allocator_init
, thrust::device(thrust::device_allocator<void>{})
, THRUST_FWD(first), THRUST_FWD(last)
-, unittest::random_integer<T>()
+, random_integer<T>()
);
DEFINE_STATEFUL_ASYNC_REDUCE_INVOKER(
reduce_async_invoker_device_on_init
@@ -173,7 +173,7 @@ DEFINE_STATEFUL_ASYNC_REDUCE_INVOKER(
// Arguments to `thrust::async::reduce`.
, thrust::device.on(stream_)
, THRUST_FWD(first), THRUST_FWD(last)
-, unittest::random_integer<T>()
+, random_integer<T>()
);
DEFINE_STATEFUL_ASYNC_REDUCE_INVOKER(
reduce_async_invoker_device_allocator_on_init
@@ -192,33 +192,33 @@ DEFINE_STATEFUL_ASYNC_REDUCE_INVOKER(
// Arguments to `thrust::async::reduce`.
, thrust::device(thrust::device_allocator<void>{}).on(stream_)
, THRUST_FWD(first), THRUST_FWD(last)
-, unittest::random_integer<T>()
+, random_integer<T>()
);
DEFINE_SYNC_REDUCE_INVOKER(
reduce_sync_invoker_init
, THRUST_FWD(first), THRUST_FWD(last)
-, unittest::random_integer<T>()
+, random_integer<T>()
);
DEFINE_ASYNC_REDUCE_INVOKER(
reduce_async_invoker_init_plus
, THRUST_FWD(first), THRUST_FWD(last)
-, unittest::random_integer<T>()
+, random_integer<T>()
, thrust::plus<T>()
);
DEFINE_ASYNC_REDUCE_INVOKER(
reduce_async_invoker_device_init_plus
, thrust::device
, THRUST_FWD(first), THRUST_FWD(last)
-, unittest::random_integer<T>()
+, random_integer<T>()
, thrust::plus<T>()
);
DEFINE_ASYNC_REDUCE_INVOKER(
reduce_async_invoker_device_allocator_init_plus
, thrust::device(thrust::device_allocator<void>{})
, THRUST_FWD(first), THRUST_FWD(last)
-, unittest::random_integer<T>()
+, random_integer<T>()
, thrust::plus<T>()
);
DEFINE_STATEFUL_ASYNC_REDUCE_INVOKER(
@@ -238,7 +238,7 @@ DEFINE_STATEFUL_ASYNC_REDUCE_INVOKER(
// Arguments to `thrust::async::reduce`.
, thrust::device.on(stream_)
, THRUST_FWD(first), THRUST_FWD(last)
-, unittest::random_integer<T>()
+, random_integer<T>()
, thrust::plus<T>()
);
DEFINE_STATEFUL_ASYNC_REDUCE_INVOKER(
@@ -258,35 +258,35 @@ DEFINE_STATEFUL_ASYNC_REDUCE_INVOKER(
// Arguments to `thrust::async::reduce`.
, thrust::device(thrust::device_allocator<void>{}).on(stream_)
, THRUST_FWD(first), THRUST_FWD(last)
-, unittest::random_integer<T>()
+, random_integer<T>()
, thrust::plus<T>()
);
DEFINE_SYNC_REDUCE_INVOKER(
reduce_sync_invoker_init_plus
, THRUST_FWD(first), THRUST_FWD(last)
-, unittest::random_integer<T>()
+, random_integer<T>()
, thrust::plus<T>()
);
DEFINE_ASYNC_REDUCE_INVOKER(
reduce_async_invoker_init_custom_plus
, THRUST_FWD(first), THRUST_FWD(last)
-, unittest::random_integer<T>()
+, random_integer<T>()
, custom_plus<T>()
);
DEFINE_ASYNC_REDUCE_INVOKER(
reduce_async_invoker_device_init_custom_plus
, thrust::device
, THRUST_FWD(first), THRUST_FWD(last)
-, unittest::random_integer<T>()
+, random_integer<T>()
, custom_plus<T>()
);
DEFINE_ASYNC_REDUCE_INVOKER(
reduce_async_invoker_device_allocator_init_custom_plus
, thrust::device(thrust::device_allocator<void>{})
, THRUST_FWD(first), THRUST_FWD(last)
-, unittest::random_integer<T>()
+, random_integer<T>()
, custom_plus<T>()
);
DEFINE_STATEFUL_ASYNC_REDUCE_INVOKER(
@@ -306,7 +306,7 @@ DEFINE_STATEFUL_ASYNC_REDUCE_INVOKER(
// Arguments to `thrust::async::reduce`.
, thrust::device.on(stream_)
, THRUST_FWD(first), THRUST_FWD(last)
-, unittest::random_integer<T>()
+, random_integer<T>()
, custom_plus<T>()
);
DEFINE_STATEFUL_ASYNC_REDUCE_INVOKER(
@@ -326,21 +326,21 @@ DEFINE_STATEFUL_ASYNC_REDUCE_INVOKER(
// Arguments to `thrust::async::reduce`.
, thrust::device(thrust::device_allocator<void>{}).on(stream_)
, THRUST_FWD(first), THRUST_FWD(last)
-, unittest::random_integer<T>()
+, random_integer<T>()
, custom_plus<T>()
);
DEFINE_SYNC_REDUCE_INVOKER(
reduce_sync_invoker_init_custom_plus
, THRUST_FWD(first), THRUST_FWD(last)
-, unittest::random_integer<T>()
+, random_integer<T>()
, custom_plus<T>()
);
///////////////////////////////////////////////////////////////////////////////
template <
- T
+ typename T
, template <typename> class AsyncReduceInvoker
, template <typename> class SyncReduceInvoker
>
@@ -384,10 +384,15 @@ void testAsyncReduce()
// This potentially runs concurrently with the copies.
auto const r0 = invoke_sync(h0.begin(), h0.end());
- auto const r1a = TEST_FUTURE_VALUE_RETRIEVAL(f0a);
- auto const r1b = TEST_FUTURE_VALUE_RETRIEVAL(f0b);
- auto const r1c = TEST_FUTURE_VALUE_RETRIEVAL(f0c);
- auto const r1d = TEST_FUTURE_VALUE_RETRIEVAL(f0d);
+ T r1a;
+ T r1b;
+ T r1c;
+ T r1d;
+
+ test_future_value_retrieval(f0a, r1a);
+ test_future_value_retrieval(f0b, r1b);
+ test_future_value_retrieval(f0c, r1c);
+ test_future_value_retrieval(f0d, r1d);
ASSERT_EQ(r0, r1a);
ASSERT_EQ(r0, r1b);
@@ -600,7 +605,7 @@ TYPED_TEST(AsyncReduceTests, TestAsyncReducePolicyAllocatorOnInitCustomPlus)
///////////////////////////////////////////////////////////////////////////////
template <
- T
+ typename T
, template <typename> class AsyncReduceInvoker
, template <typename> class SyncReduceInvoker
>
@@ -608,7 +613,7 @@ void testAsyncReduceCountingIterator()
{
constexpr std::size_t n = 15 * sizeof(T);
- ASSERT_LEQUAL(T(n), unittest::truncate_to_max_representable<T>(n));
+ ASSERT_LE(T(n), truncate_to_max_representable<T>(n));
thrust::counting_iterator<T> first(0);
thrust::counting_iterator<T> last(n);
@@ -629,10 +634,15 @@ void testAsyncReduceCountingIterator()
// This potentially runs concurrently with the copies.
auto const r0 = invoke_sync(first, last);
- auto const r1a = TEST_FUTURE_VALUE_RETRIEVAL(f0a);
- auto const r1b = TEST_FUTURE_VALUE_RETRIEVAL(f0b);
- auto const r1c = TEST_FUTURE_VALUE_RETRIEVAL(f0c);
- auto const r1d = TEST_FUTURE_VALUE_RETRIEVAL(f0d);
+ T r1a;
+ T r1b;
+ T r1c;
+ T r1d;
+
+ test_future_value_retrieval(f0a, r1a);
+ test_future_value_retrieval(f0b, r1b);
+ test_future_value_retrieval(f0c, r1c);
+ test_future_value_retrieval(f0d, r1d);
ASSERT_EQ(r0, r1a);
ASSERT_EQ(r0, r1b);
@@ -761,8 +771,11 @@ TYPED_TEST(AsyncReduceTests, TestAsyncReduceUsing)
// This potentially runs concurrently with the copies.
T const r0 = reduce(h0.begin(), h0.end());
- T const r1a = TEST_FUTURE_VALUE_RETRIEVAL(f0a);
- T const r1b = TEST_FUTURE_VALUE_RETRIEVAL(f0b);
+ T r1a;
+ T r1b;
+
+ test_future_value_retrieval(f0a, r1a);
+ test_future_value_retrieval(f0b, r1b);
ASSERT_EQ(r0, r1a);
ASSERT_EQ(r0, r1b);
@@ -804,13 +817,12 @@ TYPED_TEST(AsyncReduceTests, TestAsyncReduceAfter)
);
// Verify that double consumption of a future produces an exception.
- ASSERT_THROWS_EQUAL(
+ ASSERT_THROW(
auto x = thrust::async::reduce(
thrust::device.after(f0), d0.begin(), d0.end()
);
THRUST_UNUSED_VAR(x)
, thrust::event_error
- , thrust::event_error(thrust::event_errc::no_state)
);
ASSERT_EQ_QUIET(f0_stream, f1.stream().native_handle());
@@ -822,13 +834,12 @@ TYPED_TEST(AsyncReduceTests, TestAsyncReduceAfter)
);
// Verify that double consumption of a policy produces an exception.
- ASSERT_THROWS_EQUAL(
+ ASSERT_THROW(
auto x = thrust::async::reduce(
after_policy2, d0.begin(), d0.end()
);
THRUST_UNUSED_VAR(x)
, thrust::event_error
- , thrust::event_error(thrust::event_errc::no_state)
);
ASSERT_EQ_QUIET(f0_stream, f2.stream().native_handle());
@@ -836,7 +847,9 @@ TYPED_TEST(AsyncReduceTests, TestAsyncReduceAfter)
// This potentially runs concurrently with the copies.
T const r0 = thrust::reduce(h0.begin(), h0.end());
- T const r1 = TEST_FUTURE_VALUE_RETRIEVAL(f2);
+ T r1;
+
+ test_future_value_retrieval(f2, r1);
ASSERT_EQ(r0, r1);
}
@@ -880,13 +893,12 @@ TYPED_TEST(AsyncReduceTests, TestAsyncReduceOnThenAfter)
);
// Verify that double consumption of a future produces an exception.
- ASSERT_THROWS_EQUAL(
+ ASSERT_THROW(
auto x = thrust::async::reduce(
thrust::device.after(f0), d0.begin(), d0.end()
);
THRUST_UNUSED_VAR(x)
, thrust::event_error
- , thrust::event_error(thrust::event_errc::no_state)
);
ASSERT_EQ_QUIET(stream, f1.stream().native_handle());
@@ -898,13 +910,12 @@ TYPED_TEST(AsyncReduceTests, TestAsyncReduceOnThenAfter)
);
// Verify that double consumption of a policy produces an exception.
- ASSERT_THROWS_EQUAL(
+ ASSERT_THROW(
auto x = thrust::async::reduce(
after_policy2, d0.begin(), d0.end()
);
THRUST_UNUSED_VAR(x)
, thrust::event_error
- , thrust::event_error(thrust::event_errc::no_state)
);
ASSERT_EQ_QUIET(stream, f2.stream().native_handle());
@@ -912,7 +923,9 @@ TYPED_TEST(AsyncReduceTests, TestAsyncReduceOnThenAfter)
// This potentially runs concurrently with the copies.
T const r0 = thrust::reduce(h0.begin(), h0.end());
- T const r1 = TEST_FUTURE_VALUE_RETRIEVAL(f2);
+ T r1;
+
+ test_future_value_retrieval(f2, r1);
ASSERT_EQ(r0, r1);
@@ -967,14 +980,13 @@ TYPED_TEST(AsyncReduceTests, TestAsyncReduceAllocatorOnThenAfter)
, d0.begin(), d0.end()
);
- ASSERT_THROWS_EQUAL(
+ ASSERT_THROW(
auto x = thrust::async::reduce(
thrust::device(thrust::device_allocator<void>{}).after(f0)
, d0.begin(), d0.end()
);
THRUST_UNUSED_VAR(x)
, thrust::event_error
- , thrust::event_error(thrust::event_errc::no_state)
);
ASSERT_EQ_QUIET(stream0, f1.stream().native_handle());
@@ -984,14 +996,13 @@ TYPED_TEST(AsyncReduceTests, TestAsyncReduceAllocatorOnThenAfter)
, d0.begin(), d0.end()
);
- ASSERT_THROWS_EQUAL(
+ ASSERT_THROW(
auto x = thrust::async::reduce(
thrust::device(thrust::device_allocator<void>{}).on(stream1).after(f1)
, d0.begin(), d0.end()
);
THRUST_UNUSED_VAR(x)
, thrust::event_error
- , thrust::event_error(thrust::event_errc::no_state)
);
KNOWN_FAILURE;
@@ -1010,7 +1021,7 @@ TYPED_TEST(AsyncReduceTests, TestAsyncReduceAllocatorOnThenAfter)
thrust::hip_rocprim::throw_on_error(hipStreamDestroy(stream1));
}
}
-}*/
+}
///////////////////////////////////////////////////////////////////////////////
@@ -1065,6 +1076,7 @@ TYPED_TEST(AsyncReduceTests, TestAsyncReduceCaching)
}
}
};
+*/
///////////////////////////////////////////////////////////////////////////////
@@ -1145,10 +1157,15 @@ TYPED_TEST(AsyncReduceTests, TestAsyncCopyThenReduce)
// This potentially runs concurrently with the copies.
T const r0 = thrust::reduce(h0a.begin(), h0a.end());
- T const r1a = TEST_FUTURE_VALUE_RETRIEVAL(f1a);
- T const r1b = TEST_FUTURE_VALUE_RETRIEVAL(f1b);
- T const r1c = TEST_FUTURE_VALUE_RETRIEVAL(f1c);
- T const r1d = TEST_FUTURE_VALUE_RETRIEVAL(f1d);
+ T r1a;
+ T r1b;
+ T r1c;
+ T r1d;
+
+ test_future_value_retrieval(f1a, r1a);
+ test_future_value_retrieval(f1b, r1b);
+ test_future_value_retrieval(f1c, r1c);
+ test_future_value_retrieval(f1d, r1d);
ASSERT_EQ(r0, r1a);
ASSERT_EQ(r0, r1b);

MODIFY
test/test_utils.hpp
test/test_utils.hpp
#include <thrust/execution_policy.h>
#include <thrust/host_vector.h>
+#include <thrust/random.h>
+#include <thrust/limits.h>
#include <thrust/mr/allocator.h>
+#include <thrust/detail/event_error.h>
#include "test_seed.hpp"
#include <vector>
#include <iterator>
+#define TEST_EVENT_WAIT(e) test_event_wait(e)
+
+// for demangling the result of type_info.name()
+// with msvc, type_info.name() is already demangled
+#ifdef __GNUC__
+#include <cxxabi.h>
+#endif // __GNUC__
+
+#include <string>
+#include <cstdlib>
+
+#ifdef __GNUC__
+inline std::string demangle(const char* name)
+{
+ int status = 0;
+ char* realname = abi::__cxa_demangle(name, 0, 0, &status);
+ std::string result(realname);
+ std::free(realname);
+
+ return result;
+}
+#else
+inline std::string demangle(const char* name)
+{
+ return name;
+}
+#endif
+
+class UnitTestException
+{
+ public:
+ std::string message;
+
+ UnitTestException() {}
+ UnitTestException(const std::string& msg) : message(msg) {}
+
+ friend std::ostream& operator<<(std::ostream& os, const UnitTestException& e)
+ {
+ return os << e.message;
+ }
+
+ template <typename T>
+ UnitTestException& operator<<(const T& t)
+ {
+ std::ostringstream oss;
+ oss << t;
+ message += oss.str();
+ return *this;
+ }
+};
+
+
+class UnitTestError : public UnitTestException
+{
+ public:
+ UnitTestError() {}
+ UnitTestError(const std::string& msg) : UnitTestException(msg) {}
+};
+
+class UnitTestKnownFailure : public UnitTestException
+{
+ public:
+ UnitTestKnownFailure() {}
+ UnitTestKnownFailure(const std::string& msg) : UnitTestException(msg) {}
+};
+
+
+
+class UnitTestFailure : public UnitTestException
+{
+ public:
+ UnitTestFailure() {}
+ UnitTestFailure(const std::string& msg) : UnitTestException(msg) {}
+};
+
+template<typename T>
+ std::string type_name(void)
+{
+ return demangle(typeid(T).name());
+} // end type_name()
+
+template <typename Event>
+__host__
+void test_event_wait(Event&& e)
+{
+ ASSERT_EQ(true, e.valid_stream());
+
+ e.wait();
+ e.wait();
+
+ ASSERT_EQ(true, e.valid_stream());
+ ASSERT_EQ(true, e.ready());
+}
+
std::vector<size_t> get_sizes()
{
std::vector<size_t> sizes = {
@@ -259,3 +356,230 @@ struct key_value
key_type key;
value_type value;
};
+
+inline unsigned int hash(unsigned int a)
+{
+ a = (a+0x7ed55d16) + (a<<12);
+ a = (a^0xc761c23c) ^ (a>>19);
+ a = (a+0x165667b1) + (a<<5);
+ a = (a+0xd3a2646c) ^ (a<<9);
+ a = (a+0xfd7046c5) + (a<<3);
+ a = (a^0xb55a4f09) ^ (a>>16);
+ return a;
+}
+
+template<typename T, typename = void>
+ struct generate_random_integer;
+
+template<typename T>
+ struct generate_random_integer<T,
+ typename thrust::detail::disable_if<
+ thrust::detail::is_non_bool_arithmetic<T>::value
+ >::type
+ >
+{
+ T operator()(unsigned int i) const
+ {
+ thrust::default_random_engine rng(hash(i));
+
+ return static_cast<T>(rng());
+ }
+};
+
+template<typename T>
+ struct generate_random_integer<T,
+ typename thrust::detail::enable_if<
+ thrust::detail::is_non_bool_integral<T>::value
+ >::type
+ >
+{
+ T operator()(unsigned int i) const
+ {
+ thrust::default_random_engine rng(hash(i));
+ thrust::uniform_int_distribution<T> dist;
+
+ return static_cast<T>(dist(rng));
+ }
+};
+
+template<typename T>
+ struct generate_random_integer<T,
+ typename thrust::detail::enable_if<
+ thrust::detail::is_floating_point<T>::value
+ >::type
+ >
+{
+ T operator()(unsigned int i) const
+ {
+ T const min = std::numeric_limits<T>::min();
+ T const max = std::numeric_limits<T>::max();
+
+ thrust::default_random_engine rng(hash(i));
+ thrust::uniform_real_distribution<T> dist(min, max);
+
+ return static_cast<T>(dist(rng));
+ }
+};
+
+template<>
+ struct generate_random_integer<bool>
+{
+ bool operator()(unsigned int i) const
+ {
+ thrust::default_random_engine rng(hash(i));
+ thrust::uniform_int_distribution<unsigned int> dist(0,1);
+
+ return dist(rng) == 1;
+ }
+};
+
+
+template<typename T>
+ struct generate_random_sample
+{
+ T operator()(unsigned int i) const
+ {
+ thrust::default_random_engine rng(hash(i));
+ thrust::uniform_int_distribution<unsigned int> dist(0,20);
+
+ return static_cast<T>(dist(rng));
+ }
+};
+
+
+
+template<typename T>
+thrust::host_vector<T> random_integers(const size_t N)
+{
+ thrust::host_vector<T> vec(N);
+ thrust::transform(thrust::counting_iterator<size_t>(0),
+ thrust::counting_iterator<size_t>(N),
+ vec.begin(),
+ generate_random_integer<T>());
+
+ return vec;
+}
+
+template<typename T>
+T random_integer()
+{
+ return generate_random_integer<T>()(0);
+}
+
+template<typename T>
+thrust::host_vector<T> random_samples(const size_t N)
+{
+ thrust::host_vector<T> vec(N);
+ thrust::transform(thrust::counting_iterator<size_t>(0),
+ thrust::counting_iterator<size_t>(N),
+ vec.begin(),
+ generate_random_sample<T>());
+
+ return vec;
+}
+
+// Use this with counting_iterator to avoid generating a range larger than we
+// can represent.
+template <typename T>
+typename thrust::detail::disable_if<
+ thrust::detail::is_floating_point<T>::value
+, T
+>::type truncate_to_max_representable(std::size_t n)
+{
+ return thrust::min<std::size_t>(
+ n, static_cast<std::size_t>(thrust::numeric_limits<T>::max())
+ );
+}
+
+// TODO: This probably won't work for `half`.
+template <typename T>
+typename thrust::detail::enable_if<
+ thrust::detail::is_floating_point<T>::value
+, T
+>::type truncate_to_max_representable(std::size_t n)
+{
+ return thrust::min<T>(
+ n, thrust::numeric_limits<T>::max()
+ );
+}
+
+enum threw_status
+{
+ did_not_throw
+, threw_wrong_type
+, threw_right_type_but_wrong_value
+, threw_right_type
+};
+
+void check_assert_throws(
+ threw_status s
+, std::string const& exception_name
+, std::string const& file_name = "unknown"
+, int line_number = -1
+)
+{
+ switch (s)
+ {
+ case did_not_throw:
+ {
+ UnitTestFailure f;
+ f << "[" << file_name << ":" << line_number << "] did not throw anything";
+ throw f;
+ }
+ case threw_wrong_type:
+ {
+ UnitTestFailure f;
+ f << "[" << file_name << ":" << line_number << "] did not throw an "
+ << "object of type " << exception_name;
+ throw f;
+ }
+ case threw_right_type_but_wrong_value:
+ {
+ UnitTestFailure f;
+ f << "[" << file_name << ":" << line_number << "] threw an object of the "
+ << "correct type (" << exception_name << ") but wrong value";
+ throw f;
+ }
+ case threw_right_type:
+ break;
+ default:
+ {
+ UnitTestFailure f;
+ f << "[" << file_name << ":" << line_number << "] encountered an "
+ << "unknown error";
+ throw f;
+ }
+ }
+}
+
+template <typename Future>
+__host__
+void test_future_value_retrieval(Future&& f, decltype(f.extract()) &return_value)
+{
+ ASSERT_EQ(true, f.valid_stream());
+ ASSERT_EQ(true, f.valid_content());
+
+ auto const r0 = f.get();
+ auto const r1 = f.get();
+
+ ASSERT_EQ(true, f.ready());
+ ASSERT_EQ(true, f.valid_stream());
+ ASSERT_EQ(true, f.valid_content());
+ ASSERT_EQ(r0, r1);
+
+ auto const r2 = f.extract();
+
+ ASSERT_THROW(
+ auto x = f.extract();
+ THRUST_UNUSED_VAR(x)
+ , thrust::event_error
+ );
+
+ ASSERT_EQ(false, f.ready());
+ ASSERT_EQ(false, f.valid_stream());
+ ASSERT_EQ(false, f.valid_content());
+ ASSERT_EQ(r2, r1);
+ ASSERT_EQ(r2, r0);
+
+ return_value = r2;
+}

MODIFY
thrust/system/hip/detail/async/reduce.h
thrust/system/hip/detail/async/reduce.h
@@ -75,7 +75,7 @@ auto async_reduce_n(
using pointer
= typename thrust::detail::allocator_traits<decltype(device_alloc)>::
- rebind_traits<U>::pointer;
+ template rebind_traits<U>::pointer;
unique_eager_future_promise_pair<U, pointer> fp;
@@ -88,9 +88,9 @@ auto async_reduce_n(
, tmp_size
, first
, static_cast<U*>(nullptr)
+ , init
, n
, op
- , init
, NULL // Null stream, just for sizing.
, THRUST_HIP_DEBUG_SYNC_FLAG
)
@@ -170,10 +170,10 @@ auto async_reduce_n(
, tmp_size
, first
, ret_ptr
+ , init
, n
, op
- , init
- , fp.future.stream()
+ , fp.future.stream().native_handle()
, THRUST_HIP_DEBUG_SYNC_FLAG
)
, "after reduction launch"

Fixing async reduce for hip and enabling test.