Update future, pointer and memory resource
MODIFY
thrust/system/hip/detail/future.inl
thrust/system/hip/detail/future.inl
#include <thrust/allocate_unique.h>
#include <thrust/detail/static_assert.h>
#include <thrust/detail/execute_with_dependencies.h>
-#include <thrust/system/cuda/memory.h>
-#include <thrust/system/cuda/future.h>
-#include <thrust/system/cuda/detail/util.h>
-#include <thrust/system/cuda/detail/get_value.h>
+#include <thrust/system/hip/memory.h>
+#include <thrust/system/hip/future.h>
+#include <thrust/system/hip/detail/util.h>
+#include <thrust/system/hip/detail/get_value.h>
#include <type_traits>
#include <memory>
THRUST_BEGIN_NS
-namespace system { namespace cuda { namespace detail
+namespace system { namespace hip { namespace detail
{
///////////////////////////////////////////////////////////////////////////////
@@ -51,7 +51,7 @@ struct event_deleter final
void operator()(CUevent_st* e) const
{
if (nullptr != e)
- thrust::cuda_cub::throw_on_error(cudaEventDestroy(e));
+ thrust::hip_rocprim::throw_on_error(hipEventDestroy(e));
}
};
@@ -72,8 +72,8 @@ public:
: handle_(nullptr, event_deleter())
{
native_handle_type e;
- thrust::cuda_cub::throw_on_error(
- cudaEventCreateWithFlags(&e, cudaEventDisableTiming)
+ thrust::hip_rocprim::throw_on_error(
+ hipEventCreateWithFlags(&e, hipEventDisableTiming)
);
handle_.reset(e);
}
@@ -102,13 +102,13 @@ public:
__host__
bool ready() const
{
- cudaError_t const err = cudaEventQuery(handle_.get());
+ hipError_t const err = hipEventQuery(handle_.get());
- if (cudaErrorNotReady == err)
+ if (hipErrorNotReady == err)
return false;
// Throw on any other error.
- thrust::cuda_cub::throw_on_error(err);
+ thrust::hip_rocprim::throw_on_error(err);
return true;
}
@@ -116,7 +116,7 @@ public:
__host__
void wait() const
{
- thrust::cuda_cub::throw_on_error(cudaEventSynchronize(handle_.get()));
+ thrust::hip_rocprim::throw_on_error(hipEventSynchronize(handle_.get()));
}
__host__
@@ -140,7 +140,7 @@ struct stream_deleter final
void operator()(CUstream_st* s) const
{
if (nullptr != s)
- thrust::cuda_cub::throw_on_error(cudaStreamDestroy(s));
+ thrust::hip_rocprim::throw_on_error(hipStreamDestroy(s));
}
};
@@ -163,7 +163,7 @@ public:
{
if (cond_ && nullptr != s)
{
- thrust::cuda_cub::throw_on_error(cudaStreamDestroy(s));
+ thrust::hip_rocprim::throw_on_error(hipStreamDestroy(s));
}
}
};
@@ -185,8 +185,8 @@ public:
: handle_(nullptr, stream_conditional_deleter())
{
native_handle_type s;
- thrust::cuda_cub::throw_on_error(
- cudaStreamCreateWithFlags(&s, cudaStreamNonBlocking)
+ thrust::hip_rocprim::throw_on_error(
+ hipStreamCreateWithFlags(&s, hipStreamNonBlocking)
);
handle_.reset(s);
}
@@ -222,13 +222,13 @@ public:
__host__
bool ready() const
{
- cudaError_t const err = cudaStreamQuery(handle_.get());
+ hipError_t const err = hipStreamQuery(handle_.get());
- if (cudaErrorNotReady == err)
+ if (hipErrorNotReady == err)
return false;
// Throw on any other error.
- thrust::cuda_cub::throw_on_error(err);
+ thrust::hip_rocprim::throw_on_error(err);
return true;
}
@@ -236,16 +236,16 @@ public:
__host__
void wait() const
{
- thrust::cuda_cub::throw_on_error(
- cudaStreamSynchronize(handle_.get())
+ thrust::hip_rocprim::throw_on_error(
+ hipStreamSynchronize(handle_.get())
);
}
__host__
void depend_on(unique_event& e)
{
- thrust::cuda_cub::throw_on_error(
- cudaStreamWaitEvent(handle_.get(), e.get(), 0)
+ thrust::hip_rocprim::throw_on_error(
+ hipStreamWaitEvent(handle_.get(), e.get(), 0)
);
}
@@ -263,7 +263,7 @@ public:
__host__
void record(unique_event& e)
{
- thrust::cuda_cub::throw_on_error(cudaEventRecord(e.get(), handle_.get()));
+ thrust::hip_rocprim::throw_on_error(hipEventRecord(e.get(), handle_.get()));
}
__host__
@@ -313,24 +313,24 @@ struct acquired_stream final
// If `acquired_from` is empty, then the stream is newly created.
};
-// Precondition: `device` is the current CUDA device.
+// Precondition: `device` is the current HIP device.
template <typename X, typename Y, typename Deleter>
__host__
optional<unique_stream>
try_acquire_stream(int device, std::unique_ptr<Y, Deleter>&) noexcept;
-// Precondition: `device` is the current CUDA device.
+// Precondition: `device` is the current HIP device.
inline __host__
optional<unique_stream>
try_acquire_stream(int, unique_stream& stream) noexcept;
-// Precondition: `device` is the current CUDA device.
+// Precondition: `device` is the current HIP device.
template <typename T>
__host__
optional<unique_stream>
try_acquire_stream(int device, ready_future<T>&) noexcept;
-// Precondition: `device` is the current CUDA device.
+// Precondition: `device` is the current HIP device.
template <typename X, typename XPointer>
__host__
optional<unique_stream>
@@ -371,7 +371,7 @@ public:
template <typename X, typename XPointer>
friend __host__
optional<unique_stream>
- thrust::system::cuda::detail::try_acquire_stream(
+ thrust::system::hip::detail::try_acquire_stream(
int device, unique_eager_future<X, XPointer>& parent
) noexcept;
};
@@ -528,7 +528,7 @@ public:
>
friend __host__
unique_eager_future_promise_pair<X, XPointer>
- thrust::system::cuda::detail::depend_on(
+ thrust::system::hip::detail::depend_on(
ComputeContent&& cc, std::tuple<Dependencies...>&& deps
);
};
@@ -564,7 +564,7 @@ public:
>
friend __host__
unique_eager_future_promise_pair<X, XPointer>
- thrust::system::cuda::detail::depend_on(
+ thrust::system::hip::detail::depend_on(
ComputeContent&& cc, std::tuple<Dependencies...>&& deps
);
};
@@ -672,7 +672,7 @@ public:
template <typename X, typename XPointer>
__host__
friend optional<detail::unique_stream>
- thrust::system::cuda::detail::try_acquire_stream(
+ thrust::system::hip::detail::try_acquire_stream(
int device, unique_eager_future<X, XPointer>& parent
) noexcept;
@@ -682,7 +682,7 @@ public:
>
friend __host__
detail::unique_eager_future_promise_pair<X, XPointer>
- thrust::system::cuda::detail::depend_on(
+ thrust::system::hip::detail::depend_on(
ComputeContent&& cc, std::tuple<Dependencies...>&& deps
);
};
@@ -741,7 +741,7 @@ public:
template <typename X, typename XPointer>
__host__
friend optional<detail::unique_stream>
- thrust::system::cuda::detail::try_acquire_stream(
+ thrust::system::hip::detail::try_acquire_stream(
int device, unique_eager_future<X, XPointer>& parent
) noexcept;
@@ -751,7 +751,7 @@ public:
>
friend __host__
detail::unique_eager_future_promise_pair<X, XPointer>
- thrust::system::cuda::detail::depend_on(
+ thrust::system::hip::detail::depend_on(
ComputeContent&& cc, std::tuple<Dependencies...>&& deps
);
};
@@ -990,7 +990,7 @@ unique_eager_future_promise_pair<X, XPointer>
depend_on(ComputeContent&& cc, std::tuple<Dependencies...>&& deps)
{
int device = 0;
- thrust::cuda_cub::throw_on_error(cudaGetDevice(&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);
@@ -1021,9 +1021,8 @@ depend_on(ComputeContent&& cc, std::tuple<Dependencies...>&& deps)
///////////////////////////////////////////////////////////////////////////////
-}} // namespace system::cuda
+}} // namespace system::hip
THRUST_END_NS
#endif // THRUST_CPP_DIALECT >= 2011
-

MODIFY
thrust/system/hip/detail/pointer.inl
thrust/system/hip/detail/pointer.inl
@@ -38,6 +38,36 @@ template<typename T>
namespace hip_rocprim {
+#if THRUST_CPP_DIALECT >= 2011
+template <typename T>
+__host__ __device__
+bool operator==(decltype(nullptr), pointer<T> p)
+{
+ return nullptr == p.get();
+}
+
+template <typename T>
+__host__ __device__
+bool operator==(pointer<T> p, decltype(nullptr))
+{
+ return nullptr == p.get();
+}
+
+template <typename T>
+__host__ __device__
+bool operator!=(decltype(nullptr), pointer<T> p)
+{
+ return !(nullptr == p);
+}
+
+template <typename T>
+__host__ __device__
+bool operator!=(pointer<T> p, decltype(nullptr))
+{
+ return !(nullptr == p);
+}
+#endif
+
template <typename T>
template <typename OtherT>
__host__ __device__ reference<T> &reference<T>::operator=(

ADD
thrust/system/hip/future.h
+// Copyright (c) 2018 NVIDIA Corporation
+// Author: Bryce Adelstein Lelbach <brycelelbach@gmail.com>
+//
+// Distributed under the Boost Software License v1.0 (boost.org/LICENSE_1_0.txt)
+
+#pragma once
+
+#include <thrust/detail/config.h>
+#include <thrust/detail/cpp11_required.h>
+
+#if THRUST_CPP_DIALECT >= 2011
+
+#include <thrust/system/hip/pointer.h>
+#include <thrust/system/hip/detail/execution_policy.h>
+
+#include <thrust/future.h>
+
+THRUST_BEGIN_NS
+
+namespace system { namespace hip
+{
+
+template <typename T>
+struct ready_future;
+
+template <typename T, typename Pointer = pointer<T>>
+struct unique_eager_future;
+
+}} // namespace system::cuda
+
+namespace hip
+{
+
+template <typename T>
+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
+
+template <typename T, typename Pointer, typename DerivedPolicy>
+__host__ __device__
+thrust::system::hip::unique_eager_future<T, Pointer>
+unique_eager_future_type(thrust::hip_rocprim::execution_policy<DerivedPolicy> const&);
+
+THRUST_END_NS
+
+#include <thrust/system/hip/detail/future.inl>
+
+#endif // THRUST_CPP_DIALECT >= 2011

MODIFY
thrust/system/hip/memory_resource.h
thrust/system/hip/memory_resource.h
#include <thrust/memory/detail/host_system_resource.h>
-namespace thrust
-{
+THRUST_BEGIN_NS
+
namespace system
{
namespace hip
@@ -103,4 +103,5 @@ typedef detail::pinned_memory_resource universal_host_pinned_memory_resource;
} // end hip
} // end system
-} // end thrust
+
+THRUST_END_NS

MODIFY
thrust/system/hip/pointer.h
thrust/system/hip/pointer.h
@@ -187,6 +187,25 @@ public:
}
}; // struct pointer
+#if THRUST_CPP_DIALECT >= 2011
+template <typename T>
+__host__ __device__
+bool operator!=(decltype(nullptr), pointer<T>);
+
+template <typename T>
+__host__ __device__
+bool operator!=(pointer<T>, decltype(nullptr));
+
+template <typename T>
+__host__ __device__
+bool operator==(decltype(nullptr), pointer<T>);
+
+template <typename T>
+__host__ __device__
+bool operator==(pointer<T>, decltype(nullptr));
+#endif
+
+
/*! \p reference is a wrapped reference to an object stored in memory available to the \p hip system.
* \p reference is the type of the result of dereferencing a \p hip::pointer.
*

Update future, pointer and memory resource