Fix reduce and extend workaround
MODIFY
thrust/detail/config/compiler.h
thrust/detail/config/compiler.h
// XXX we should move the definition of THRUST_DEPRECATED out of this logic
#if defined(_MSC_VER)
#define THRUST_HOST_COMPILER THRUST_HOST_COMPILER_MSVC
+#define THRUST_DEPRECATED __declspec(deprecated)
#elif defined(__clang__)
#define THRUST_HOST_COMPILER THRUST_HOST_COMPILER_CLANG
#define THRUST_DEPRECATED __attribute__ ((deprecated))
#define THRUST_GCC_VERSION (__GNUC__ * 10000 + __GNUC_MINOR__ * 100 + __GNUC_PATCHLEVEL__)
#else
#define THRUST_HOST_COMPILER THRUST_HOST_COMPILER_UNKNOWN
+#define THRUST_DEPRECATED
#endif // THRUST_HOST_COMPILER
// figure out which device compiler we're using
@@ -122,21 +124,8 @@ __THRUST_DISABLE_MSVC_WARNING_BEGIN(4800)
#define __THRUST_DISABLE_MSVC_FORCING_VALUE_TO_BOOL_END \
__THRUST_DISABLE_MSVC_WARNING_END(4800)
-// figure out which host compiler we're using
-// XXX we should move the definition of THRUST_DEPRECATED out of this logic
-#if THRUST_HOST_COMPILER == THRUST_HOST_COMPILER_MSVC
- #define THRUST_DEPRECATED __declspec(deprecated)
-#elif THRUST_HOST_COMPILER == THRUST_HOST_COMPILER_CLANG
- #define THRUST_DEPRECATED __attribute__((deprecated))
-#elif THRUST_HOST_COMPILER == THRUST_HOST_COMPILER_GCC
- #define THRUST_DEPRECATED __attribute__((deprecated))
-#else
- #define THRUST_DEPRECATED
-#endif
-
#if __cplusplus >= 201103L
#define THRUST_NOEXCEPT noexcept
#else
#define THRUST_NOEXCEPT throw()
#endif
-

MODIFY
thrust/system/hip/detail/assign_value.h
thrust/system/hip/detail/assign_value.h
#pragma once
#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+#include <thrust/system/hip/config.h>
#include <thrust/detail/config.h>
#include <thrust/detail/raw_pointer_cast.h>
-#include <thrust/system/hip/config.h>
#include <thrust/system/hip/detail/copy.h>
#include <thrust/system/hip/detail/execution_policy.h>

MODIFY
thrust/system/hip/detail/for_each.h
thrust/system/hip/detail/for_each.h
#include <iterator>
#include <thrust/system/hip/config.h>
+#include <thrust/system/hip/detail/util.h>
+#include <thrust/system/hip/detail/parallel_for.h>
#include <thrust/detail/function.h>
#include <thrust/distance.h>
-#include <thrust/system/hip/detail/parallel_for.h>
-#include <thrust/system/hip/detail/util.h>
BEGIN_NS_THRUST
@@ -74,7 +74,9 @@ for_each_n(execution_policy<Derived>& policy, Input first, Size count, UnaryOp o
typedef thrust::detail::wrapped_function<UnaryOp, void> wrapped_t;
wrapped_t wrapped_op(op);
- hip_rocprim::parallel_for(policy, for_each_f<Input, wrapped_t>(first, wrapped_op), count);
+ hip_rocprim::parallel_for(policy,
+ for_each_f<Input, wrapped_t>(first, wrapped_op),
+ count);
return first + count;
}

MODIFY
thrust/system/hip/detail/internal/copy_cross_system.h
thrust/system/hip/detail/internal/copy_cross_system.h
@@ -105,27 +105,16 @@ namespace __copy
// non-trivial H->D copy
template <class H, class D, class InputIt, class Size, class OutputIt>
- OutputIt __host__ /* WORKAROUND */ __device__
- cross_system_copy_n(thrust::cpp::execution_policy<H>& host_s,
+ OutputIt __host__
+ cross_system_copy_n_hd_nt(thrust::cpp::execution_policy<H>& host_s,
thrust::hip_rocprim::execution_policy<D>& device_s,
InputIt first,
Size num_items,
- OutputIt result,
- thrust::detail::false_type) // non-trivial copy
+ OutputIt result)
{
// get type of the input data
typedef typename thrust::iterator_value<InputIt>::type InputTy;
- // WORKAROUND
-#if defined(THRUST_HIP_DEVICE_CODE)
- THRUST_UNUSED_VAR(host_s);
- THRUST_UNUSED_VAR(device_s);
- THRUST_UNUSED_VAR(first);
- THRUST_UNUSED_VAR(num_items);
- THRUST_HIP_PRESERVE_KERNELS_WORKAROUND(
- (hip_rocprim::copy_n<D, InputTy*, Size, OutputIt>));
- return result;
-#else
// copy input data into host temp storage
InputIt last = first;
thrust::advance(last, num_items);
@@ -147,39 +136,53 @@ namespace __copy
// device->device copy
OutputIt ret = hip_rocprim::copy_n(device_s, d_in_ptr.data(), num_items, result);
+
return ret;
-#endif
}
-#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, class InputIt, class Size, class OutputIt>
- OutputIt __host__ /* WORKAROUND */ __device__
- cross_system_copy_n(thrust::hip_rocprim::execution_policy<D>& device_s,
- thrust::cpp::execution_policy<H>& host_s,
+ // non-trivial H->D copy
+ template <class H, class D, class InputIt, class Size, class OutputIt>
+ OutputIt __host__ __device__
+ cross_system_copy_n(thrust::cpp::execution_policy<H>& host_s,
+ thrust::hip_rocprim::execution_policy<D>& device_s,
InputIt first,
Size num_items,
OutputIt result,
thrust::detail::false_type) // non-trivial copy
-
{
// get type of the input data
typedef typename thrust::iterator_value<InputIt>::type InputTy;
- // WORKAROUND
+ THRUST_HIP_PRESERVE_KERNELS_WORKAROUND(
+ (hip_rocprim::__copy::cross_system_copy_n_hd_nt<H, D, InputIt, Size, OutputIt>)
+ );
+
#if defined(THRUST_HIP_DEVICE_CODE)
- THRUST_UNUSED_VAR(device_s);
THRUST_UNUSED_VAR(host_s);
+ THRUST_UNUSED_VAR(device_s);
THRUST_UNUSED_VAR(first);
THRUST_UNUSED_VAR(num_items);
- THRUST_HIP_PRESERVE_KERNELS_WORKAROUND(
- (hip_rocprim::uninitialized_copy_n<D, InputIt, Size, InputTy*>)
- );
return result;
#else
+ return cross_system_copy_n_hd_nt(host_s, device_s, first, num_items, result);
+#endif
+ }
+
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+
+// non-trivial copy
+template <class D, class H, class InputIt, class Size, class OutputIt>
+OutputIt __host__
+cross_system_copy_n_dh_nt(thrust::hip_rocprim::execution_policy<D>& device_s,
+ thrust::cpp::execution_policy<H>& host_s,
+ InputIt first,
+ Size num_items,
+ OutputIt result)
+{
+ // get type of the input data
+ typedef typename thrust::iterator_value<InputIt>::type InputTy;
+
// allocate device temp storage
thrust::detail::temporary_array<InputTy, D> d_in_ptr(device_s, num_items);
@@ -198,6 +201,37 @@ namespace __copy
OutputIt ret = thrust::copy_n(host_s, temp.data(), num_items, result);
return ret;
+}
+
+// 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, class InputIt, class Size, class OutputIt>
+OutputIt __host__ __device__
+cross_system_copy_n(thrust::hip_rocprim::execution_policy<D>& device_s,
+ thrust::cpp::execution_policy<H>& host_s,
+ InputIt first,
+ Size num_items,
+ OutputIt result,
+ thrust::detail::false_type) // non-trivial copy
+
+{
+ // get type of the input data
+ typedef typename thrust::iterator_value<InputIt>::type InputTy;
+
+ THRUST_HIP_PRESERVE_KERNELS_WORKAROUND(
+ (hip_rocprim::__copy::cross_system_copy_n_dh_nt<D, H, InputIt, Size, OutputIt>)
+ );
+
+#if defined(THRUST_HIP_DEVICE_CODE)
+ THRUST_UNUSED_VAR(device_s);
+ THRUST_UNUSED_VAR(host_s);
+ THRUST_UNUSED_VAR(first);
+ THRUST_UNUSED_VAR(num_items);
+
+ return result;
+#else
+ return cross_system_copy_n_dh_nt(device_s, host_s, first, num_items, result);
#endif
}
#endif

MODIFY
thrust/system/hip/detail/parallel_for.h
thrust/system/hip/detail/parallel_for.h
#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
#include <thrust/system/hip/config.h>
+#include <thrust/system/hip/detail/util.h>
#include <thrust/detail/type_traits/result_of_adaptable_function.h>
#include <thrust/system/hip/detail/par_to_seq.h>
-#include <thrust/system/hip/detail/util.h>
BEGIN_NS_THRUST
namespace hip_rocprim

MODIFY
thrust/system/hip/detail/reduce.h
thrust/system/hip/detail/reduce.h
#pragma once
#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
-#include <thrust/detail/config.h>
+#include <thrust/system/hip/config.h>
#include <thrust/detail/cstdint.h>
#include <thrust/detail/temporary_array.h>
-#include <thrust/detail/alignment.h>
-#include <thrust/detail/minmax.h>
+#include <thrust/system/hip/detail/util.h>
#include <thrust/detail/raw_reference_cast.h>
#include <thrust/detail/type_traits/iterator/is_output_iterator.h>
-#include <thrust/system/hip/detail/get_value.h>
#include <thrust/system/hip/detail/par_to_seq.h>
-#include <thrust/system/hip/detail/util.h>
+#include <thrust/system/hip/detail/get_value.h>
+#include <thrust/functional.h>
#include <thrust/device_vector.h>
+#include <thrust/detail/minmax.h>
#include <thrust/distance.h>
-#include <thrust/functional.h>
+#include <thrust/detail/alignment.h>
// rocprim include
@@ -51,7 +51,10 @@ BEGIN_NS_THRUST
// forward declare generic reduce
// to circumvent circular dependency
-template <typename DerivedPolicy, typename InputIterator, typename T, typename BinaryFunction>
+template <typename DerivedPolicy,
+ typename InputIterator,
+ typename T,
+ typename BinaryFunction>
__host__ __device__
T reduce(const thrust::detail::execution_policy_base<DerivedPolicy>& exec,
InputIterator first,
@@ -63,7 +66,11 @@ namespace hip_rocprim
{
namespace __reduce
{
- template <typename Derived, typename InputIt, typename Size, typename T, typename BinaryOp>
+ template <typename Derived,
+ typename InputIt,
+ typename Size,
+ typename T,
+ typename BinaryOp>
THRUST_HIP_RUNTIME_FUNCTION
T reduce(execution_policy<Derived>& policy, InputIt first, Size num_items, T init, BinaryOp binary_op)
{
@@ -95,7 +102,7 @@ namespace __reduce
void *ptr = static_cast<void*>(tmp.data().get());
T* d_ret_ptr = reinterpret_cast<T*>(
- reinterpret_cast<T*>(ptr) + temp_storage_bytes);
+ reinterpret_cast<char*>(ptr) + temp_storage_bytes);
hip_rocprim::throw_on_error(rocprim::reduce(ptr,
temp_storage_bytes,

MODIFY
thrust/system/hip/detail/uninitialized_copy.h
thrust/system/hip/detail/uninitialized_copy.h
#include <iterator>
#include <thrust/distance.h>
#include <thrust/system/hip/detail/execution_policy.h>
-#include <thrust/system/hip/detail/parallel_for.h>
#include <thrust/system/hip/detail/util.h>
+#include <thrust/system/hip/detail/parallel_for.h>
BEGIN_NS_THRUST
namespace hip_rocprim

Fix reduce and extend workaround