Fix detail/async files to use the correct hip and rocprim modules
MODIFY
thrust/system/hip/detail/async/customization.h
thrust/system/hip/detail/async/customization.h
*
******************************************************************************/
-// TODO: Move into system::cuda
+// TODO: Move into system::hip
#pragma once
#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_NVCC
-#include <thrust/system/cuda/config.h>
+#include <thrust/system/hip/config.h>
#include <thrust/detail/type_deduction.h>
#include <thrust/detail/cstdint.h>
#include <thrust/detail/execute_with_allocator.h>
-#include <thrust/system/cuda/memory_resource.h>
+#include <thrust/system/hip/memory_resource.h>
#include <thrust/memory/detail/host_system_resource.h>
#include <thrust/mr/allocator.h>
#include <thrust/mr/disjoint_sync_pool.h>
THRUST_BEGIN_NS
-namespace system { namespace cuda { namespace detail
+namespace system { namespace hip { namespace detail
{
using default_async_host_resource =
@@ -71,7 +71,7 @@ THRUST_DECLTYPE_RETURNS(
using default_async_device_resource =
thrust::mr::disjoint_synchronized_pool_resource<
- thrust::system::cuda::memory_resource
+ thrust::system::hip::memory_resource
, thrust::mr::new_delete_resource
>;
@@ -95,7 +95,7 @@ THRUST_DECLTYPE_RETURNS(exec.get_allocator())
using default_async_universal_host_pinned_resource =
thrust::mr::synchronized_pool_resource<
- thrust::system::cuda::universal_host_pinned_memory_resource
+ thrust::system::hip::universal_host_pinned_memory_resource
>;
template <typename DerivedPolicy>
@@ -108,7 +108,7 @@ THRUST_DECLTYPE_RETURNS(
>{}
)
-}}} // namespace system::cuda::detail
+}}} // namespace system::hip::detail
THRUST_END_NS

MODIFY
thrust/system/hip/detail/async/for_each.h
thrust/system/hip/detail/async/for_each.h
*
******************************************************************************/
-// TODO: Move into system::cuda
+// TODO: Move into system::hip
#pragma once
#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_NVCC
-#include <thrust/system/cuda/config.h>
+#include <thrust/system/hip/config.h>
-#include <thrust/system/cuda/detail/async/customization.h>
-#include <thrust/system/cuda/detail/parallel_for.h>
-#include <thrust/system/cuda/future.h>
+#include <thrust/system/hip/detail/async/customization.h>
+#include <thrust/system/hip/detail/parallel_for.h>
+#include <thrust/system/hip/future.h>
#include <thrust/iterator/iterator_traits.h>
#include <thrust/distance.h>
THRUST_BEGIN_NS
-namespace system { namespace cuda { namespace detail
+namespace system { namespace hip { namespace detail
{
template <typename ForwardIt, typename UnaryFunction>
@@ -89,9 +89,9 @@ auto async_for_each_n(
// Set up stream with dependencies.
- cudaStream_t const user_raw_stream = thrust::cuda_cub::stream(policy);
+ hipStream_t const user_raw_stream = thrust::hip_rocprim::stream(policy);
- if (thrust::cuda_cub::default_stream() != user_raw_stream)
+ if (thrust::hip_rocprim::default_stream() != user_raw_stream)
{
fp = depend_on<void, pointer>(
nullptr
@@ -114,8 +114,8 @@ auto async_for_each_n(
std::move(first), std::move(f)
);
- thrust::cuda_cub::throw_on_error(
- thrust::cuda_cub::__parallel_for::parallel_for(
+ thrust::hip_rocprim::throw_on_error(
+ thrust::hip_rocprim::__parallel_for::parallel_for(
n, std::move(wrapped), fp.future.stream()
)
, "after for_each launch"
@@ -124,9 +124,9 @@ auto async_for_each_n(
return std::move(fp.future);
}
-}}} // namespace system::cuda::detail
+}}} // namespace system::hip::detail
-namespace cuda_cub
+namespace hip_rocprim
{
// ADL entry point.
@@ -142,12 +142,12 @@ auto async_for_each(
UnaryFunction&& f
)
THRUST_DECLTYPE_RETURNS(
- thrust::system::cuda::detail::async_for_each_n(
+ thrust::system::hip::detail::async_for_each_n(
policy, first, thrust::distance(first, last), THRUST_FWD(f)
)
);
-} // cuda_cub
+} // hip_rocprim
THRUST_END_NS

MODIFY
thrust/system/hip/detail/async/reduce.h
thrust/system/hip/detail/async/reduce.h
// TODO: Optimize for thrust::plus
-// TODO: Move into system::cuda
+// TODO: Move into system::hip
#pragma once
#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_NVCC
-#include <thrust/system/cuda/config.h>
+#include <thrust/system/hip/config.h>
-#include <thrust/system/cuda/detail/async/customization.h>
-#include <thrust/system/cuda/detail/reduce.h>
-#include <thrust/system/cuda/future.h>
+#include <thrust/system/hip/detail/async/customization.h>
+#include <thrust/system/hip/detail/reduce.h>
+#include <thrust/system/hip/future.h>
#include <thrust/iterator/iterator_traits.h>
#include <thrust/distance.h>
THRUST_BEGIN_NS
-namespace system { namespace cuda { namespace detail
+namespace system { namespace hip { namespace detail
{
template <
@@ -83,8 +83,8 @@ auto async_reduce_n(
// Determine temporary device storage requirements.
size_t tmp_size = 0;
- thrust::cuda_cub::throw_on_error(
- thrust::cuda_cub::cub::DeviceReduce::Reduce(
+ thrust::hip_rocprim::throw_on_error(
+ thrust::hip_rocprim::rocprim::DeviceReduce::Reduce(
NULL
, tmp_size
, first
@@ -105,7 +105,7 @@ auto async_reduce_n(
);
// The array was dynamically allocated, so we assume that it's suitably
- // aligned for any type of data. `malloc`/`cudaMalloc`/`new`/`std::allocator`
+ // aligned for any type of data. `malloc`/`hipMalloc`/`new`/`std::allocator`
// make this guarantee.
auto const content_ptr = content.get();
T* const ret_ptr = thrust::detail::aligned_reinterpret_cast<T*>(
@@ -117,9 +117,9 @@ auto async_reduce_n(
// Set up stream with dependencies.
- cudaStream_t const user_raw_stream = thrust::cuda_cub::stream(policy);
+ hipStream_t const user_raw_stream = thrust::hip_rocprim::stream(policy);
- if (thrust::cuda_cub::default_stream() != user_raw_stream)
+ if (thrust::hip_rocprim::default_stream() != user_raw_stream)
{
fp = depend_on<T, pointer>(
[] (decltype(content) const& c)
@@ -165,8 +165,8 @@ auto async_reduce_n(
// Run reduction.
- thrust::cuda_cub::throw_on_error(
- thrust::cuda_cub::cub::DeviceReduce::Reduce(
+ thrust::hip_rocprim::throw_on_error(
+ thrust::hip_rocprim::rocprim::DeviceReduce::Reduce(
tmp_ptr
, tmp_size
, first
@@ -183,9 +183,9 @@ auto async_reduce_n(
return std::move(fp.future);
}
-}}} // namespace system::cuda::detail
+}}} // namespace system::hip::detail
-namespace cuda_cub
+namespace hip_rocprim
{
// ADL entry point.
@@ -202,12 +202,12 @@ auto async_reduce(
BinaryOp op
)
THRUST_DECLTYPE_RETURNS(
- thrust::system::cuda::detail::async_reduce_n(
+ thrust::system::hip::detail::async_reduce_n(
policy, first, thrust::distance(first, last), init, op
)
)
-} // cuda_cub
+} // hip_rocprim
THRUST_END_NS

MODIFY
thrust/system/hip/detail/async/sort.h
thrust/system/hip/detail/async/sort.h
*
******************************************************************************/
-// TODO: Move into system::cuda
+// TODO: Move into system::hip
#pragma once
#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_NVCC
-#include <thrust/system/cuda/config.h>
+#include <thrust/system/hip/config.h>
-#include <thrust/system/cuda/detail/async/customization.h>
-#include <thrust/system/cuda/detail/sort.h>
+#include <thrust/system/hip/detail/async/customization.h>
+#include <thrust/system/hip/detail/sort.h>
#include <thrust/detail/alignment.h>
-#include <thrust/system/cuda/future.h>
+#include <thrust/system/hip/future.h>
#include <thrust/type_traits/is_trivially_relocatable.h>
#include <thrust/type_traits/is_contiguous_iterator.h>
#include <thrust/type_traits/logical_metafunctions.h>
THRUST_BEGIN_NS
-namespace system { namespace cuda { namespace detail
+namespace system { namespace hip { namespace detail
{
// Non-ContiguousIterator iterators
@@ -134,8 +134,8 @@ auto async_stable_sort_n(
// Determine temporary device storage requirements.
size_t tmp_size = 0;
- thrust::cuda_cub::throw_on_error(
- thrust::cuda_cub::__merge_sort::doit_step<
+ thrust::hip_rocprim::throw_on_error(
+ thrust::hip_rocprim::__merge_sort::doit_step<
/* Sort items? */ std::false_type, /* Stable? */ std::true_type
>(
nullptr
@@ -157,7 +157,7 @@ auto async_stable_sort_n(
);
// The array was dynamically allocated, so we assume that it's suitably
- // aligned for any type of data. `malloc`/`cudaMalloc`/`new`/`std::allocator`
+ // aligned for any type of data. `malloc`/`hipMalloc`/`new`/`std::allocator`
// make this guarantee.
auto const content_ptr = content.get();
@@ -167,9 +167,9 @@ auto async_stable_sort_n(
// Set up stream with dependencies.
- cudaStream_t const user_raw_stream = thrust::cuda_cub::stream(policy);
+ hipStream_t const user_raw_stream = thrust::hip_rocprim::stream(policy);
- if (thrust::cuda_cub::default_stream() != user_raw_stream)
+ if (thrust::hip_rocprim::default_stream() != user_raw_stream)
{
fp = depend_on<void, pointer>(
nullptr
@@ -191,8 +191,8 @@ auto async_stable_sort_n(
// Run merge sort.
- thrust::cuda_cub::throw_on_error(
- thrust::cuda_cub::__merge_sort::doit_step<
+ thrust::hip_rocprim::throw_on_error(
+ thrust::hip_rocprim::__merge_sort::doit_step<
/* Sort items? */ std::false_type, /* Stable? */ std::true_type
>(
tmp_ptr
@@ -213,7 +213,7 @@ auto async_stable_sort_n(
// ContiguousIterator iterators
// Scalar value type
// thrust::greater<>
-// TODO (hack up CUB)
+// TODO (hack up rocprim)
// ContiguousIterator iterators
// Scalar value type
@@ -254,15 +254,15 @@ auto async_stable_sort_n(
unique_eager_future_promise_pair<void, pointer> fp;
- thrust::cuda_cub::cub::DoubleBuffer<T> keys(
+ thrust::hip_rocprim::rocprim::DoubleBuffer<T> keys(
raw_pointer_cast(&*first), nullptr
);
// Determine temporary device storage requirements.
size_t tmp_size = 0;
- thrust::cuda_cub::throw_on_error(
- thrust::cuda_cub::cub::DeviceRadixSort::SortKeys(
+ thrust::hip_rocprim::throw_on_error(
+ thrust::hip_rocprim::rocprim::DeviceRadixSort::SortKeys(
nullptr
, tmp_size
, keys
@@ -286,7 +286,7 @@ auto async_stable_sort_n(
);
// The array was dynamically allocated, so we assume that it's suitably
- // aligned for any type of data. `malloc`/`cudaMalloc`/`new`/`std::allocator`
+ // aligned for any type of data. `malloc`/`hipMalloc`/`new`/`std::allocator`
// make this guarantee.
auto const content_ptr = content.get();
@@ -300,9 +300,9 @@ auto async_stable_sort_n(
// Set up stream with dependencies.
- cudaStream_t const user_raw_stream = thrust::cuda_cub::stream(policy);
+ hipStream_t const user_raw_stream = thrust::hip_rocprim::stream(policy);
- if (thrust::cuda_cub::default_stream() != user_raw_stream)
+ if (thrust::hip_rocprim::default_stream() != user_raw_stream)
{
fp = depend_on<void, pointer>(
nullptr
@@ -324,8 +324,8 @@ auto async_stable_sort_n(
// Run radix sort.
- thrust::cuda_cub::throw_on_error(
- thrust::cuda_cub::cub::DeviceRadixSort::SortKeys(
+ thrust::hip_rocprim::throw_on_error(
+ thrust::hip_rocprim::rocprim::DeviceRadixSort::SortKeys(
tmp_ptr
, tmp_size
, keys
@@ -341,12 +341,12 @@ auto async_stable_sort_n(
if (0 != keys.selector)
{
// TODO: Temporary hack.
- thrust::cuda_cub::throw_on_error(
- cudaMemcpyAsync(
+ thrust::hip_rocprim::throw_on_error(
+ hipMemcpyAsync(
reinterpret_cast<T*>(keys.d_buffers[0])
, reinterpret_cast<T*>(keys.d_buffers[1])
, sizeof(T) * n
- , cudaMemcpyDeviceToDevice
+ , hipMemcpyDeviceToDevice
, fp.future.stream()
)
, "radix sort copy back"
@@ -356,9 +356,9 @@ auto async_stable_sort_n(
return std::move(fp.future);
}
-}}} // namespace system::cuda::detail
+}}} // namespace system::hip::detail
-namespace cuda_cub
+namespace hip_rocprim
{
// ADL entry point.
@@ -374,12 +374,12 @@ auto async_stable_sort(
StrictWeakOrdering comp
)
THRUST_DECLTYPE_RETURNS(
- thrust::system::cuda::detail::async_stable_sort_n(
+ thrust::system::hip::detail::async_stable_sort_n(
policy, first, thrust::distance(first, last), comp
)
);
-} // cuda_cub
+} // hip_rocprim
THRUST_END_NS

MODIFY
thrust/system/hip/detail/async/transform.h
thrust/system/hip/detail/async/transform.h
*
******************************************************************************/
-// TODO: Move into system::cuda
+// TODO: Move into system::hip
#pragma once
#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_NVCC
-#include <thrust/system/cuda/config.h>
+#include <thrust/system/hip/config.h>
-#include <thrust/system/cuda/detail/async/customization.h>
-#include <thrust/system/cuda/detail/parallel_for.h>
-#include <thrust/system/cuda/future.h>
+#include <thrust/system/hip/detail/async/customization.h>
+#include <thrust/system/hip/detail/parallel_for.h>
+#include <thrust/system/hip/future.h>
#include <thrust/iterator/iterator_traits.h>
#include <thrust/distance.h>
THRUST_BEGIN_NS
-namespace system { namespace cuda { namespace detail
+namespace system { namespace hip { namespace detail
{
template <typename ForwardIt, typename OutputIt, typename UnaryOperation>
@@ -108,9 +108,9 @@ auto async_transform_n(
// Set up stream with dependencies.
- cudaStream_t const user_raw_stream = thrust::cuda_cub::stream(policy);
+ hipStream_t const user_raw_stream = thrust::hip_rocprim::stream(policy);
- if (thrust::cuda_cub::default_stream() != user_raw_stream)
+ if (thrust::hip_rocprim::default_stream() != user_raw_stream)
{
fp = depend_on<return_type, return_pointer>(
[] (decltype(content) const& c)
@@ -138,8 +138,8 @@ auto async_transform_n(
std::move(first), std::move(output), std::move(op)
);
- thrust::cuda_cub::throw_on_error(
- thrust::cuda_cub::__parallel_for::parallel_for(
+ thrust::hip_rocprim::throw_on_error(
+ thrust::hip_rocprim::__parallel_for::parallel_for(
n, std::move(wrapped), fp.future.stream()
)
, "after transform launch"
@@ -148,9 +148,9 @@ auto async_transform_n(
return std::move(fp.future);
}
-}}} // namespace system::cuda::detail
+}}} // namespace system::hip::detail
-namespace cuda_cub
+namespace hip_rocprim
{
// ADL entry point.
@@ -168,12 +168,12 @@ auto async_transform(
UnaryOperation&& op
)
THRUST_DECLTYPE_RETURNS(
- thrust::system::cuda::detail::async_transform_n(
+ thrust::system::hip::detail::async_transform_n(
policy, first, thrust::distance(first, last), output, THRUST_FWD(op)
)
);
-} // cuda_cub
+} // hip_rocprim
THRUST_END_NS

Fix detail/async files to use the correct hip and rocprim modules