Fixing multiple issues regarding namespace usage and functions not being defined in the correct place.
MODIFY
thrust/system/hip/config.h
thrust/system/hip/config.h
#ifndef __HIP_DEVICE_COMPILE__
#define __THRUST_HAS_HIPRT__ 1
#define THRUST_HIP_RUNTIME_FUNCTION __host__ __device__ __forceinline__
+ #define THRUST_RUNTIME_FUNCTION THRUST_HIP_RUNTIME_FUNCTION
#else
#define __THRUST_HAS_HIPRT__ 0
#define THRUST_HIP_RUNTIME_FUNCTION __host__ __forceinline__
+ #define THRUST_RUNTIME_FUNCTION THRUST_HIP_RUNTIME_FUNCTION
#endif
#else
#define __THRUST_HAS_HIPRT__ 0
#define THRUST_HIP_RUNTIME_FUNCTION __host__ __forceinline__
+ #define THRUST_RUNTIME_FUNCTION THRUST_HIP_RUNTIME_FUNCTION
#endif
// TODO: These paremeters should be tuned for NAVI.

MODIFY
thrust/system/hip/detail/async/reduce.h
thrust/system/hip/detail/async/reduce.h
@@ -84,7 +84,7 @@ auto async_reduce_n(
size_t tmp_size = 0;
thrust::hip_rocprim::throw_on_error(
- thrust::hip_rocprim::rocprim::DeviceReduce::Reduce(
+ rocprim::reduce(
NULL
, tmp_size
, first
@@ -93,7 +93,7 @@ auto async_reduce_n(
, op
, init
, NULL // Null stream, just for sizing.
- , THRUST_DEBUG_SYNC_FLAG
+ , THRUST_HIP_DEBUG_SYNC_FLAG
)
, "after reduction sizing"
);
@@ -117,7 +117,7 @@ auto async_reduce_n(
// Set up stream with dependencies.
- hipStream_t const user_raw_stream = thrust::hip_rocprim::stream(policy);
+ hipStream_t user_raw_stream = thrust::hip_rocprim::stream(policy);
if (thrust::hip_rocprim::default_stream() != user_raw_stream)
{
@@ -133,7 +133,7 @@ auto async_reduce_n(
, std::tuple_cat(
std::make_tuple(
std::move(content)
- , unique_stream(nonowning, user_raw_stream)
+ , unique_stream(nonowning, &user_raw_stream)
)
, extract_dependencies(
std::move(thrust::detail::derived_cast(policy))
@@ -166,7 +166,7 @@ auto async_reduce_n(
// Run reduction.
thrust::hip_rocprim::throw_on_error(
- thrust::hip_rocprim::rocprim::DeviceReduce::Reduce(
+ rocprim::reduce(
tmp_ptr
, tmp_size
, first
@@ -175,7 +175,7 @@ auto async_reduce_n(
, op
, init
, fp.future.stream()
- , THRUST_DEBUG_SYNC_FLAG
+ , THRUST_HIP_DEBUG_SYNC_FLAG
)
, "after reduction launch"
);

MODIFY
thrust/system/hip/detail/count.h
thrust/system/hip/detail/count.h
@@ -48,7 +48,7 @@ count_if(execution_policy<Derived>& policy,
typedef typename iterator_traits<InputIt>::difference_type size_type;
typedef transform_input_iterator_t<size_type, InputIt, UnaryPred> flag_iterator_t;
- return hip_rocprim::reduce_n(policy,
+ return reduce_n(policy,
flag_iterator_t(first, unary_pred),
thrust::distance(first, last),
size_type(0),
@@ -62,7 +62,7 @@ count(execution_policy<Derived>& policy,
InputIt last,
Value const& value)
{
- return hip_rocprim::count_if(
+ return count_if(
policy, first, last, thrust::detail::equal_to_value<Value>(value)
);
}

MODIFY
thrust/system/hip/detail/future.inl
thrust/system/hip/detail/future.inl
@@ -48,10 +48,10 @@ constexpr nonowning_t nonowning{};
struct event_deleter final
{
__host__
- void operator()(CUevent_st* e) const
+ void operator()(hipEvent_t* e) const
{
if (nullptr != e)
- thrust::hip_rocprim::throw_on_error(hipEventDestroy(e));
+ thrust::hip_rocprim::throw_on_error(hipEventDestroy(*e), "");
}
};
@@ -59,10 +59,10 @@ struct event_deleter final
struct unique_event final
{
- using native_handle_type = CUevent_st*;
+ using native_handle_type = hipEvent_t*;
private:
- std::unique_ptr<CUevent_st, event_deleter> handle_;
+ std::unique_ptr<hipEvent_t, event_deleter> handle_;
public:
/// \brief Create a new stream and construct a handle to it. When the handle
@@ -71,9 +71,10 @@ public:
unique_event()
: handle_(nullptr, event_deleter())
{
- native_handle_type e;
+ native_handle_type e = nullptr;
thrust::hip_rocprim::throw_on_error(
- hipEventCreateWithFlags(&e, hipEventDisableTiming)
+ hipEventCreateWithFlags(e, hipEventDisableTiming),
+ ""
);
handle_.reset(e);
}
@@ -102,13 +103,13 @@ public:
__host__
bool ready() const
{
- hipError_t const err = hipEventQuery(handle_.get());
+ hipError_t const err = hipEventQuery(*handle_.get());
if (hipErrorNotReady == err)
return false;
// Throw on any other error.
- thrust::hip_rocprim::throw_on_error(err);
+ thrust::hip_rocprim::throw_on_error(err, "");
return true;
}
@@ -116,7 +117,7 @@ public:
__host__
void wait() const
{
- thrust::hip_rocprim::throw_on_error(hipEventSynchronize(handle_.get()));
+ thrust::hip_rocprim::throw_on_error(hipEventSynchronize(*handle_.get()), "");
}
__host__
@@ -137,10 +138,10 @@ public:
struct stream_deleter final
{
__host__
- void operator()(CUstream_st* s) const
+ void operator()(hipStream_t* s) const
{
if (nullptr != s)
- thrust::hip_rocprim::throw_on_error(hipStreamDestroy(s));
+ thrust::hip_rocprim::throw_on_error(hipStreamDestroy(*s), "");
}
};
@@ -159,11 +160,11 @@ public:
: cond_(false) {}
__host__
- void operator()(CUstream_st* s) const
+ void operator()(hipStream_t* s) const
{
if (cond_ && nullptr != s)
{
- thrust::hip_rocprim::throw_on_error(hipStreamDestroy(s));
+ thrust::hip_rocprim::throw_on_error(hipStreamDestroy(*s), "");
}
}
};
@@ -172,10 +173,10 @@ public:
struct unique_stream final
{
- using native_handle_type = CUstream_st*;
+ using native_handle_type = hipStream_t*;
private:
- std::unique_ptr<CUstream_st, stream_conditional_deleter> handle_;
+ std::unique_ptr<hipStream_t, stream_conditional_deleter> handle_;
public:
/// \brief Create a new stream and construct a handle to it. When the handle
@@ -184,9 +185,10 @@ public:
unique_stream()
: handle_(nullptr, stream_conditional_deleter())
{
- native_handle_type s;
+ native_handle_type s = nullptr;
thrust::hip_rocprim::throw_on_error(
- hipStreamCreateWithFlags(&s, hipStreamNonBlocking)
+ hipStreamCreateWithFlags(s, hipStreamNonBlocking),
+ ""
);
handle_.reset(s);
}
@@ -222,13 +224,13 @@ public:
__host__
bool ready() const
{
- hipError_t const err = hipStreamQuery(handle_.get());
+ hipError_t const err = hipStreamQuery(*handle_.get());
if (hipErrorNotReady == err)
return false;
// Throw on any other error.
- thrust::hip_rocprim::throw_on_error(err);
+ thrust::hip_rocprim::throw_on_error(err, "");
return true;
}
@@ -237,7 +239,8 @@ public:
void wait() const
{
thrust::hip_rocprim::throw_on_error(
- hipStreamSynchronize(handle_.get())
+ hipStreamSynchronize(*handle_.get()),
+ ""
);
}
@@ -245,7 +248,8 @@ public:
void depend_on(unique_event& e)
{
thrust::hip_rocprim::throw_on_error(
- hipStreamWaitEvent(handle_.get(), e.get(), 0)
+ hipStreamWaitEvent(*handle_.get(), *e.get(), 0),
+ ""
);
}
@@ -263,7 +267,7 @@ public:
__host__
void record(unique_event& e)
{
- thrust::hip_rocprim::throw_on_error(hipEventRecord(e.get(), handle_.get()));
+ thrust::hip_rocprim::throw_on_error(hipEventRecord(*e.get(), *handle_.get()), "");
}
__host__
@@ -990,7 +994,7 @@ unique_eager_future_promise_pair<X, XPointer>
depend_on(ComputeContent&& cc, std::tuple<Dependencies...>&& deps)
{
int device = 0;
- thrust::hip_rocprim::throw_on_error(hipGetDevice(&device));
+ thrust::hip_rocprim::throw_on_error(hipGetDevice(&device), "");
// First, either steal a stream from one of our children or make a new one.
auto as = acquire_stream(device, deps);

MODIFY
thrust/system/hip/detail/inner_product.h
thrust/system/hip/detail/inner_product.h
@@ -61,7 +61,7 @@ inner_product(execution_policy<Derived>& policy,
typedef transform_pair_of_input_iterators_t<T, InputIt1, InputIt2, ProductOp>
binop_iterator_t;
- return hip_rocprim::reduce_n(
+ return reduce_n(
policy, binop_iterator_t(first1, first2, product_op), num_items, init, reduce_op
);
}
@@ -74,7 +74,7 @@ inner_product(execution_policy<Derived>& policy,
InputIt2 first2,
T init)
{
- return hip_rocprim::inner_product(
+ return inner_product(
policy, first1, last1, first2, init, plus<T>(), multiplies<T>()
);
}

MODIFY
thrust/system/hip/detail/transform_reduce.h
thrust/system/hip/detail/transform_reduce.h
@@ -49,7 +49,7 @@ transform_reduce(execution_policy<Derived>& policy,
size_type num_items = static_cast<size_type>(thrust::distance(first, last));
typedef transform_input_iterator_t<T, InputIt, TransformOp> transformed_iterator_t;
- return hip_rocprim::reduce_n(
+ return reduce_n(
policy, transformed_iterator_t(first, transform_op), num_items, init, reduce_op
);
}

MODIFY
thrust/system/hip/future.h
thrust/system/hip/future.h
@@ -26,7 +26,7 @@ struct ready_future;
template <typename T, typename Pointer = pointer<T>>
struct unique_eager_future;
-}} // namespace system::cuda
+}} // namespace system::hip
namespace hip
{
@@ -37,7 +37,7 @@ using ready_future = thrust::system::hip::ready_future<T>;
template <typename T, typename Pointer = thrust::system::hip::pointer<T>>
using unique_eager_future = thrust::system::hip::unique_eager_future<T, Pointer>;
-} // namespace cuda
+} // namespace hip
template <typename T, typename Pointer, typename DerivedPolicy>
__host__ __device__

MODIFY
thrust/system/hip/pointer.h
thrust/system/hip/pointer.h
@@ -46,12 +46,13 @@ public:
typedef typename ptr::reference reference;
};
+namespace thrust
+{
+
// specialize thrust::iterator_traits to avoid problems with the name of
// pointer's constructor shadowing its nested pointer type
// do this before pointer is defined so the specialization is correctly
// used inside the definition
-namespace thrust
-{
template <typename Element>
struct iterator_traits<thrust::hip_rocprim::pointer<Element> >

Fixing multiple issues regarding namespace usage and functions not being defined in the correct place.