HIP backend: port copying operations
MODIFY
thrust/system/hip/detail/assign_value.h
thrust/system/hip/detail/assign_value.h
#include <thrust/system/hip/config.h>
#include <thrust/system/hip/detail/execution_policy.h>
#include <thrust/detail/raw_pointer_cast.h>
-// #include <thrust/system/hip/detail/copy.h>
+#include <thrust/system/hip/detail/copy.h>
BEGIN_NS_THRUST
@@ -32,25 +32,11 @@ template<typename DerivedPolicy, typename Pointer1, typename Pointer2>
inline __host__ __device__
void assign_value(thrust::hip::execution_policy<DerivedPolicy> &exec, Pointer1 dst, Pointer2 src)
{
- // XXX war nvbugs/881631
- struct war_nvbugs_881631
- {
- __host__ inline static void host_path(thrust::hip::execution_policy<DerivedPolicy> &exec, Pointer1 dst, Pointer2 src)
- {
- (void) exec; (void) dst; (void) src;
- // hip_rocprim::copy(exec, src, src + 1, dst);
- }
-
- __device__ inline static void device_path(thrust::hip::execution_policy<DerivedPolicy> &, Pointer1 dst, Pointer2 src)
- {
+#ifdef __HIP_DEVICE_COMPILE__
+ (void) exec;
*thrust::raw_pointer_cast(dst) = *thrust::raw_pointer_cast(src);
- }
- };
-
-#ifndef __HIP_DEVICE_COMPILE__
- war_nvbugs_881631::host_path(exec,dst,src);
#else
- war_nvbugs_881631::device_path(exec,dst,src);
+ hip_rocprim::copy(exec, src, src + 1, dst);
#endif // __HIP_DEVICE_COMPILE__
} // end assign_value()
@@ -59,37 +45,16 @@ template<typename System1, typename System2, typename Pointer1, typename Pointer
inline __host__ __device__
void assign_value(cross_system<System1,System2> &systems, Pointer1 dst, Pointer2 src)
{
- // XXX war nvbugs/881631
- struct war_nvbugs_881631
- {
- __host__ inline static void host_path(cross_system<System1,System2> &systems, Pointer1 dst, Pointer2 src)
- {
- // rotate the systems so that they are ordered the same as (src, dst)
- // for the call to thrust::copy
- cross_system<System2,System1> rotated_systems = systems.rotate();
- (void) rotated_systems; (void) dst; (void) src;
- // hip_rocprim::copy(rotated_systems, src, src + 1, dst);
- }
-
- __device__ inline static void device_path(cross_system<System1,System2> &, Pointer1 dst, Pointer2 src)
- {
- // XXX forward the true hip::execution_policy inside systems here
- // instead of materializing a tag
+#ifdef __HIP_DEVICE_COMPILE__
+ (void) systems;
thrust::hip::tag hip_tag;
thrust::hip_rocprim::assign_value(hip_tag, dst, src);
- }
- };
-
-#if __HIP_DEVICE_COMPILE__
- war_nvbugs_881631::device_path(systems,dst,src);
#else
- war_nvbugs_881631::host_path(systems,dst,src);
+ cross_system<System2,System1> rotated_systems = systems.rotate();
+ hip_rocprim::copy(rotated_systems, src, src + 1, dst);
#endif
} // end assign_value()
-
-
-
} // end hip_rocprim
END_NS_THRUST
#endif

MODIFY
thrust/system/hip/detail/copy.h
thrust/system/hip/detail/copy.h
#pragma once
-#include <thrust/system/cuda/config.h>
-#include <thrust/system/cuda/detail/execution_policy.h>
-#include <thrust/system/cuda/detail/cross_system.h>
+#include <thrust/system/hip/config.h>
+#include <thrust/system/hip/detail/execution_policy.h>
+#include <thrust/system/hip/detail/cross_system.h>
BEGIN_NS_THRUST
@@ -47,9 +47,9 @@ copy_n(const thrust::detail::execution_policy_base<DerivedPolicy> &exec,
Size n,
OutputIt result);
-namespace cuda_cub {
+namespace hip_rocprim {
-// D->D copy requires NVCC compiler
+// D->D copy requires HCC compiler
template <class System,
class InputIterator,
class OutputIterator>
@@ -90,21 +90,21 @@ copy_n(cross_system<System1, System2> systems,
Size n,
OutputIterator result);
-} // namespace cuda_
+} // namespace hip_rocprim
END_NS_THRUST
-#include <thrust/system/cuda/detail/internal/copy_device_to_device.h>
-#include <thrust/system/cuda/detail/internal/copy_cross_system.h>
-#include <thrust/system/cuda/detail/par_to_seq.h>
+#include <thrust/system/hip/detail/internal/copy_device_to_device.h>
+#include <thrust/system/hip/detail/internal/copy_cross_system.h>
+#include <thrust/system/hip/detail/par_to_seq.h>
BEGIN_NS_THRUST
-namespace cuda_cub {
+namespace hip_rocprim {
-#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_NVCC
-// D->D copy requires NVCC compiler
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+// D->D copy requires HCC compiler
__thrust_exec_check_disable__
template <class System,
@@ -117,20 +117,24 @@ copy(execution_policy<System> &system,
OutputIterator result)
{
OutputIterator ret = result;
- if (__THRUST_HAS_CUDART__)
+ {
+ auto ptr = __copy::device_to_device<
+ System, InputIterator, OutputIterator
+ >;
+ (void) ptr;
+ }
+#if __THRUST_HAS_HIPRT__
{
ret = __copy::device_to_device(system, first, last, result);
}
- else
+#else
{
-#if !__THRUST_HAS_CUDART__
ret = thrust::copy(cvt_to_seq(derived_cast(system)),
first,
last,
result);
-#endif
}
-
+#endif
return ret;
} // end copy()
@@ -146,17 +150,21 @@ copy_n(execution_policy<System> &system,
OutputIterator result)
{
OutputIterator ret = result;
- if (__THRUST_HAS_CUDART__)
+ {
+ auto ptr = __copy::device_to_device<
+ System, InputIterator, OutputIterator
+ >;
+ (void) ptr;
+ }
+#if __THRUST_HAS_HIPRT__
{
ret = __copy::device_to_device(system, first, first + n, result);
}
- else
+#else
{
-#if !__THRUST_HAS_CUDART__
ret = thrust::copy_n(cvt_to_seq(derived_cast(system)), first, n, result);
-#endif
}
-
+#endif
return ret;
} // end copy_n()
#endif
@@ -189,7 +197,7 @@ copy_n(cross_system<System1, System2> systems,
} // end copy_n()
-} // namespace cuda_cub
+} // namespace hip_rocprim
END_NS_THRUST
#include <thrust/memory.h>

MODIFY
thrust/system/hip/detail/cross_system.h
thrust/system/hip/detail/cross_system.h
******************************************************************************/
#pragma once
-#include <thrust/system/cuda/detail/cross_system.h>
+#include <thrust/system/cpp/detail/execution_policy.h>
+#include <thrust/system/hip/detail/execution_policy.h>
BEGIN_NS_THRUST
namespace hip_rocprim {
- using ::thrust::cuda_cub::cross_system;
- using ::thrust::cuda_cub::select_system;
+ template <class Sys1, class Sys2>
+ struct cross_system : thrust::execution_policy<cross_system<Sys1, Sys2> >
+ {
+ typedef thrust::execution_policy<Sys1> policy1;
+ typedef thrust::execution_policy<Sys2> policy2;
-// STREAMHPC Remove that code below or stop using cross_system, select_system from CUDA backend
+ policy1 &sys1;
+ policy2 &sys2;
-// template <class Sys1, class Sys2>
-// struct cross_system : thrust::execution_policy<cross_system<Sys1, Sys2> >
-// {
-// typedef thrust::execution_policy<Sys1> policy1;
-// typedef thrust::execution_policy<Sys2> policy2;
+ inline __host__ __device__
+ cross_system(policy1 &sys1, policy2 &sys2) : sys1(sys1), sys2(sys2) {}
-// policy1 &sys1;
-// policy2 &sys2;
+ __host__ __device__ inline cross_system<Sys2, Sys1>
+ rotate() const
+ {
+ return cross_system<Sys2, Sys1>(sys2, sys1);
+ }
+ };
-// inline __host__ __device__
-// cross_system(policy1 &sys1, policy2 &sys2) : sys1(sys1), sys2(sys2) {}
+ // host interop: (device,host)
+ template <class Sys1, class Sys2>
+ __host__ __device__ inline cross_system<Sys1, Sys2>
+ select_system(execution_policy<Sys1> const & sys1,
+ thrust::cpp::execution_policy<Sys2> const &sys2)
+ {
+ thrust::execution_policy<Sys1> & non_const_sys1 = const_cast<execution_policy<Sys1> &>(sys1);
+ thrust::cpp::execution_policy<Sys2> &non_const_sys2 = const_cast<thrust::cpp::execution_policy<Sys2> &>(sys2);
+ return cross_system<Sys1, Sys2>(non_const_sys1, non_const_sys2);
+ }
-// __host__ __device__ inline cross_system<Sys2, Sys1>
-// rotate() const
-// {
-// return cross_system<Sys2, Sys1>(sys2, sys1);
-// }
-// };
-
-// // host interop: (device,host)
-// template <class Sys1, class Sys2>
-// __host__ __device__ inline cross_system<Sys1, Sys2>
-// select_system(execution_policy<Sys1> const & sys1,
-// thrust::cpp::execution_policy<Sys2> const &sys2)
-// {
-// thrust::execution_policy<Sys1> & non_const_sys1 = const_cast<execution_policy<Sys1> &>(sys1);
-// thrust::cpp::execution_policy<Sys2> &non_const_sys2 = const_cast<thrust::cpp::execution_policy<Sys2> &>(sys2);
-// return cross_system<Sys1, Sys2>(non_const_sys1, non_const_sys2);
-// }
-
-// // host interop: (host,device)
-// template <class Sys1, class Sys2>
-// __host__ __device__ inline cross_system<Sys1, Sys2>
-// select_system(const thrust::cpp::execution_policy<Sys1> &sys1, execution_policy<Sys2> &sys2)
-// {
-// thrust::cpp::execution_policy<Sys1> &non_const_sys1 = const_cast<thrust::cpp::execution_policy<Sys1> &>(sys1);
-// thrust::execution_policy<Sys2> & non_const_sys2 = const_cast<execution_policy<Sys2> &>(sys2);
-// return cross_system<Sys1, Sys2>(non_const_sys1, non_const_sys2);
-// }
+ // host interop: (host,device)
+ template <class Sys1, class Sys2>
+ __host__ __device__ inline cross_system<Sys1, Sys2>
+ select_system(const thrust::cpp::execution_policy<Sys1> &sys1, execution_policy<Sys2> &sys2)
+ {
+ thrust::cpp::execution_policy<Sys1> &non_const_sys1 = const_cast<thrust::cpp::execution_policy<Sys1> &>(sys1);
+ thrust::execution_policy<Sys2> & non_const_sys2 = const_cast<execution_policy<Sys2> &>(sys2);
+ return cross_system<Sys1, Sys2>(non_const_sys1, non_const_sys2);
+ }
} // namespace hip_rocprim
END_NS_THRUST

MODIFY
thrust/system/hip/detail/internal/copy_cross_system.h
thrust/system/hip/detail/internal/copy_cross_system.h
// XXX
// this file must not be included on its own, ever,
-// but must be part of include in thrust/system/cuda/detail/copy.h
+// but must be part of include in thrust/system/hip/detail/copy.h
-#include <thrust/system/cuda/config.h>
+#include <thrust/system/hip/config.h>
#include <thrust/detail/dispatch/is_trivial_copy.h>
#include <thrust/distance.h>
#include <thrust/advance.h>
#include <thrust/detail/raw_pointer_cast.h>
-#include <thrust/system/cuda/detail/uninitialized_copy.h>
+#include <thrust/system/hip/detail/uninitialized_copy.h>
#include <thrust/detail/temporary_buffer.h>
BEGIN_NS_THRUST
-namespace cuda_cub {
+namespace hip_rocprim {
namespace __copy {
@@ -50,38 +50,42 @@ namespace __copy {
class D,
class T,
class Size>
- THRUST_HOST_FUNCTION void
+ THRUST_HIP_HOST_FUNCTION void
trivial_device_copy(thrust::cpp::execution_policy<H>& ,
- thrust::cuda_cub::execution_policy<D>& device_s,
+ thrust::hip_rocprim::execution_policy<D>& device_s,
T* dst,
T const* src,
Size count)
{
- cudaError status;
- status = cuda_cub::trivial_copy_to_device(dst,
+ hipError_t status;
+ status = hip_rocprim::trivial_copy_to_device(
+ dst,
src,
count,
- cuda_cub::stream(device_s));
- cuda_cub::throw_on_error(status, "__copy::trivial_device_copy H->D: failed");
+ hip_rocprim::stream(device_s)
+ );
+ hip_rocprim::throw_on_error(status, "__copy::trivial_device_copy H->D: failed");
}
template <class D,
class H,
class T,
class Size>
- THRUST_HOST_FUNCTION void
- trivial_device_copy(thrust::cuda_cub::execution_policy<D>& device_s,
+ THRUST_HIP_HOST_FUNCTION void
+ trivial_device_copy(thrust::hip_rocprim::execution_policy<D>& device_s,
thrust::cpp::execution_policy<H>& ,
T* dst,
T const* src,
Size count)
{
- cudaError status;
- status = cuda_cub::trivial_copy_from_device(dst,
+ hipError_t status;
+ status = hip_rocprim::trivial_copy_from_device(
+ dst,
src,
count,
- cuda_cub::stream(device_s));
- cuda_cub::throw_on_error(status, "trivial_device_copy D->H failed");
+ hip_rocprim::stream(device_s)
+ );
+ hip_rocprim::throw_on_error(status, "trivial_device_copy D->H failed");
}
template <class System1,
@@ -117,7 +121,7 @@ namespace __copy {
class OutputIt>
OutputIt __host__
cross_system_copy_n(thrust::cpp::execution_policy<H>& host_s,
- thrust::cuda_cub::execution_policy<D>& device_s,
+ thrust::hip_rocprim::execution_policy<D>& device_s,
InputIt first,
Size num_items,
OutputIt result,
@@ -144,22 +148,24 @@ namespace __copy {
// allocate device temporary storage
- cudaError status;
+ hipError_t status;
InputTy* d_in_ptr = thrust::raw_pointer_cast(
thrust::get_temporary_buffer<InputTy>(
device_s, sizeof(InputTy) * num_items)
.first);
// trivial copy data from host to device
- status = cuda_cub::trivial_copy_to_device(d_in_ptr,
+ status = hip_rocprim::trivial_copy_to_device(
+ d_in_ptr,
temp,
num_items,
- cuda_cub::stream(device_s));
- cuda_cub::throw_on_error(status, "__copy:: H->D: failed");
+ hip_rocprim::stream(device_s)
+ );
+ hip_rocprim::throw_on_error(status, "__copy:: H->D: failed");
// device->device copy
- OutputIt ret = cuda_cub::copy_n(device_s, d_in_ptr,num_items, result);
+ OutputIt ret = hip_rocprim::copy_n(device_s, d_in_ptr,num_items, result);
// free device temporary storage
thrust::return_temporary_buffer(host_s, temp);
@@ -168,9 +174,9 @@ namespace __copy {
return ret;
}
-#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_NVCC
- // non-trivial copy D->H, only supported with NVCC compiler
- // because copy ctor must have __device__ annotations ,which is nvcc-only
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+ // non-trivial copy D->H, only supported with HCC compiler
+ // because copy ctor must have __device__ annotations, which is hcc-only
// feature
template <class D,
class H,
@@ -178,7 +184,7 @@ namespace __copy {
class Size,
class OutputIt>
OutputIt __host__
- cross_system_copy_n(thrust::cuda_cub::execution_policy<D>& device_s,
+ cross_system_copy_n(thrust::hip_rocprim::execution_policy<D>& device_s,
thrust::cpp::execution_policy<H>& host_s,
InputIt first,
Size num_items,
@@ -190,14 +196,14 @@ namespace __copy {
typedef typename thrust::iterator_value<InputIt>::type InputTy;
// allocate device temp storage
- cudaError status;
+ hipError_t status;
InputTy* d_in_ptr = thrust::raw_pointer_cast(
thrust::get_temporary_buffer<InputTy>(
device_s, sizeof(InputTy) * num_items).first);
// uninitialize copy into temp device storage
- cuda_cub::uninitialized_copy_n(device_s, first,num_items, d_in_ptr);
+ hip_rocprim::uninitialized_copy_n(device_s, first,num_items, d_in_ptr);
// allocate host temp storage
// thrust::detail::temporary_array<InputTy,H> temp(0, host_s, num_items);
@@ -205,11 +211,13 @@ namespace __copy {
thrust::get_temporary_buffer<InputTy>(host_s,num_items).first);
// trivial copy from device to host
- status = cuda_cub::trivial_copy_from_device(temp,
+ status = hip_rocprim::trivial_copy_from_device(
+ temp,
d_in_ptr,
num_items,
- cuda_cub::stream(device_s));
- cuda_cub::throw_on_error(status, "__copy:: D->H: failed");
+ hip_rocprim::stream(device_s)
+ );
+ hip_rocprim::throw_on_error(status, "__copy:: D->H: failed");
// copy host->host
@@ -269,5 +277,5 @@ namespace __copy {
} // namespace __copy
-} // namespace cuda_cub
+} // namespace hip_rocprim
END_NS_THRUST

MODIFY
thrust/system/hip/detail/internal/copy_device_to_device.h
thrust/system/hip/detail/internal/copy_device_to_device.h
#pragma once
-#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_NVCC
-#include <thrust/system/cuda/config.h>
-#include <thrust/system/cuda/detail/execution_policy.h>
-#include <thrust/system/cuda/detail/transform.h>
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+#include <thrust/system/hip/config.h>
+#include <thrust/system/hip/detail/execution_policy.h>
+#include <thrust/system/hip/detail/transform.h>
#include <thrust/functional.h>
+#include <thrust/distance.h>
BEGIN_NS_THRUST
-namespace cuda_cub {
+namespace hip_rocprim {
namespace __copy {
template <class Derived,
class InputIt,
class OutputIt>
- OutputIt CUB_RUNTIME_FUNCTION
+ OutputIt THRUST_HIP_HOST_FUNCTION
device_to_device(execution_policy<Derived>& policy,
InputIt first,
InputIt last,
OutputIt result)
{
typedef typename thrust::iterator_traits<InputIt>::value_type InputTy;
- return cuda_cub::transform(policy,
+ return hip_rocprim::transform(
+ policy,
first,
last,
result,
- thrust::identity<InputTy>());
+ thrust::identity<InputTy>()
+ );
}
} // namespace __copy
-} // namespace cuda_cub
+} // namespace hip_rocprim
END_NS_THRUST
#endif

MODIFY
thrust/system/hip/detail/uninitialized_copy.h
thrust/system/hip/detail/uninitialized_copy.h
#pragma once
-#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_NVCC
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
#include <iterator>
#include <thrust/distance.h>
-#include <thrust/system/cuda/detail/execution_policy.h>
-#include <thrust/system/cuda/detail/util.h>
-#include <thrust/system/cuda/detail/parallel_for.h>
+#include <thrust/system/hip/detail/execution_policy.h>
+#include <thrust/system/hip/detail/util.h>
+#include <thrust/system/hip/detail/parallel_for.h>
BEGIN_NS_THRUST
-namespace cuda_cub {
+namespace hip_rocprim {
namespace __uninitialized_copy {
@@ -49,17 +49,18 @@ namespace __uninitialized_copy {
typedef typename iterator_traits<InputIt>::value_type InputType;
typedef typename iterator_traits<OutputIt>::value_type OutputType;
- THRUST_FUNCTION
+ THRUST_HIP_FUNCTION
functor(InputIt input_, OutputIt output_)
: input(input_), output(output_) {}
template<class Size>
- void THRUST_DEVICE_FUNCTION operator()(Size idx)
+ void THRUST_HIP_DEVICE_FUNCTION operator()(Size idx)
{
InputType const &in = raw_reference_cast(input[idx]);
OutputType & out = raw_reference_cast(output[idx]);
-#if defined(__CUDA__) && defined(__clang__)
+// STREAMHPC Check if this is a problem for HCC too
+#if defined(__HCC__) && defined(__clang__)
// XXX unsafe, but clang is seemngly unable to call in-place new
out = in;
#else
@@ -82,7 +83,7 @@ uninitialized_copy_n(execution_policy<Derived> &policy,
{
typedef __uninitialized_copy::functor<InputIt,OutputIt> functor_t;
- cuda_cub::parallel_for(policy,
+ hip_rocprim::parallel_for(policy,
functor_t(first, result),
count);
return result + count;
@@ -97,13 +98,13 @@ uninitialized_copy(execution_policy<Derived>& policy,
InputIt last,
OutputIt result)
{
- return cuda_cub::uninitialized_copy_n(policy,
+ return hip_rocprim::uninitialized_copy_n(policy,
first,
thrust::distance(first, last),
result);
}
-} // namespace cuda_
+} // namespace hip_rocprim
END_NS_THRUST
#endif

MODIFY
thrust/system/hip/execution_policy.h
thrust/system/hip/execution_policy.h
#include <thrust/detail/config.h>
#include <thrust/system/hip/detail/execution_policy.h>
-// #include <thrust/system/hip/detail/par.h>
+#include <thrust/system/hip/detail/par.h>
// // pass
// // ----------------
-// #include <thrust/system/cuda/detail/adjacent_difference.h>
-// #include <thrust/system/cuda/detail/copy.h>
-// #include <thrust/system/cuda/detail/copy_if.h>
-// #include <thrust/system/cuda/detail/count.h>
-// #include <thrust/system/cuda/detail/equal.h>
-// #include <thrust/system/cuda/detail/extrema.h>
-// #include <thrust/system/cuda/detail/fill.h>
-// #include <thrust/system/cuda/detail/find.h>
-// #include <thrust/system/cuda/detail/for_each.h>
-// #include <thrust/system/cuda/detail/gather.h>
-// #include <thrust/system/cuda/detail/generate.h>
-// #include <thrust/system/cuda/detail/inner_product.h>
-// #include <thrust/system/cuda/detail/mismatch.h>
-// #include <thrust/system/cuda/detail/partition.h>
-// #include <thrust/system/cuda/detail/reduce_by_key.h>
-// #include <thrust/system/cuda/detail/remove.h>
-// #include <thrust/system/cuda/detail/replace.h>
-// #include <thrust/system/cuda/detail/reverse.h>
-// #include <thrust/system/cuda/detail/scatter.h>
-// #include <thrust/system/cuda/detail/swap_ranges.h>
-// #include <thrust/system/cuda/detail/tabulate.h>
-// #include <thrust/system/cuda/detail/transform.h>
-// #include <thrust/system/cuda/detail/transform_reduce.h>
-// #include <thrust/system/cuda/detail/transform_scan.h>
-// #include <thrust/system/cuda/detail/uninitialized_copy.h>
-// #include <thrust/system/cuda/detail/uninitialized_fill.h>
-// #include <thrust/system/cuda/detail/unique.h>
-// #include <thrust/system/cuda/detail/unique_by_key.h>
+// #include <thrust/system/hip/detail/adjacent_difference.h>
+#include <thrust/system/hip/detail/copy.h>
+// #include <thrust/system/hip/detail/copy_if.h>
+// #include <thrust/system/hip/detail/count.h>
+// #include <thrust/system/hip/detail/equal.h>
+// #include <thrust/system/hip/detail/extrema.h>
+#include <thrust/system/hip/detail/fill.h>
+// #include <thrust/system/hip/detail/find.h>
+#include <thrust/system/hip/detail/for_each.h>
+// #include <thrust/system/hip/detail/gather.h>
+#include <thrust/system/hip/detail/generate.h>
+// #include <thrust/system/hip/detail/inner_product.h>
+// #include <thrust/system/hip/detail/mismatch.h>
+// #include <thrust/system/hip/detail/partition.h>
+// #include <thrust/system/hip/detail/reduce_by_key.h>
+// #include <thrust/system/hip/detail/remove.h>
+#include <thrust/system/hip/detail/replace.h>
+// #include <thrust/system/hip/detail/reverse.h>
+// #include <thrust/system/hip/detail/scatter.h>
+#include <thrust/system/hip/detail/swap_ranges.h>
+// #include <thrust/system/hip/detail/tabulate.h>
+#include <thrust/system/hip/detail/transform.h>
+// #include <thrust/system/hip/detail/transform_reduce.h>
+// #include <thrust/system/hip/detail/transform_scan.h>
+#include <thrust/system/hip/detail/uninitialized_copy.h>
+#include <thrust/system/hip/detail/uninitialized_fill.h>
+// #include <thrust/system/hip/detail/unique.h>
+// #include <thrust/system/hip/detail/unique_by_key.h>
// // fail
// // ----------------
// // fails with mixed types
-// #include <thrust/system/cuda/detail/reduce.h>
+// #include <thrust/system/hip/detail/reduce.h>
// // mixed types are not compiling, commented in testing/scan.cu
-// #include <thrust/system/cuda/detail/scan.h>
+// #include <thrust/system/hip/detail/scan.h>
// // stubs passed
// // ----------------
-// #include <thrust/system/cuda/detail/binary_search.h>
-// #include <thrust/system/cuda/detail/merge.h>
-// #include <thrust/system/cuda/detail/scan_by_key.h>
-// #include <thrust/system/cuda/detail/set_operations.h>
-// #include <thrust/system/cuda/detail/sort.h>
+// #include <thrust/system/hip/detail/binary_search.h>
+// #include <thrust/system/hip/detail/merge.h>
+// #include <thrust/system/hip/detail/scan_by_key.h>
+// #include <thrust/system/hip/detail/set_operations.h>
+// #include <thrust/system/hip/detail/sort.h>
// // work in progress

HIP backend: port copying operations