Update copy workaround and code formatting
MODIFY
test/test_partition.cpp
test/test_partition.cpp
@@ -1933,7 +1933,6 @@ TEST(PartitionTests,TestStablePartitionDevice)
SCOPED_TRACE(testing::Message() << "with seed= " << seed);
thrust::host_vector<int> h_data = get_random_data<int>(size, 0, size, seed);
-
thrust::device_vector<int> d_data = h_data;
thrust::stable_partition(h_data.begin(), h_data.end(),is_even<int>());

MODIFY
thrust/system/hip/detail/copy.h
thrust/system/hip/detail/copy.h
@@ -100,20 +100,41 @@ copy(execution_policy<System>& system,
InputIterator last,
OutputIterator result)
{
- OutputIterator ret = result;
+ // struct workaround is required for HIP-clang
+ // THRUST_HIP_PRESERVE_KERNELS_WORKAROUND is required for HCC
+ struct workaround
+ {
+ __host__
+ static OutputIterator par(
+ execution_policy<System>& system,
+ InputIterator first,
+ InputIterator last,
+ OutputIterator result)
+ {
+ #if __HCC__ && __HIP_DEVICE_COMPILE__
THRUST_HIP_PRESERVE_KERNELS_WORKAROUND(
(__copy::device_to_device<System, InputIterator, OutputIterator>)
);
-#if __THRUST_HAS_HIPRT__
- {
- ret = __copy::device_to_device(system, first, last, result);
- }
#else
+ return __copy::device_to_device(system, first, last, result);
+ #endif
+ }
+ __device__
+ static OutputIterator seq(
+ execution_policy<System>& system,
+ InputIterator first,
+ InputIterator last,
+ OutputIterator result)
{
- ret = thrust::copy(cvt_to_seq(derived_cast(system)), first, last, result);
+ return thrust::copy(cvt_to_seq(derived_cast(system)), first, last, result);
}
+ };
+
+#if __THRUST_HAS_HIPRT__
+ return workaround::par(policy, first, last, result);
+#else
+ return workaround::seq(policy, first, last, result);
#endif
- return ret;
} // end copy()
__thrust_exec_check_disable__ template <class System,
@@ -126,20 +147,35 @@ copy_n(execution_policy<System>& system,
Size n,
OutputIterator result)
{
- OutputIterator ret = result;
+ // struct workaround is required for HIP-clang
+ // THRUST_HIP_PRESERVE_KERNELS_WORKAROUND is required for HCC
+ struct workaround
+ {
+ __host__
+ static OutputIterator par(
+ execution_policy<System>& system,
+ InputIterator first,
+ Size n,
+ OutputIterator result)
+ {
+ #if __HCC__ && __HIP_DEVICE_COMPILE__
THRUST_HIP_PRESERVE_KERNELS_WORKAROUND(
(__copy::device_to_device<System, InputIterator, OutputIterator>)
);
-#if __THRUST_HAS_HIPRT__
- {
- ret = __copy::device_to_device(system, first, first + n, result);
- }
#else
+ return __copy::device_to_device(system, first, first + n, result);
+ #endif
+ }
+ __device__
+ static OutputIterator seq(
+ execution_policy<System>& system,
+ InputIterator first,
+ Size n,
+ OutputIterator result)
{
- ret = thrust::copy_n(cvt_to_seq(derived_cast(system)), first, n, result);
+ return thrust::copy_n(cvt_to_seq(derived_cast(system)), first, n, result);
}
-#endif
- return ret;
+ };
} // end copy_n()
#endif

MODIFY
thrust/system/hip/detail/merge.h
thrust/system/hip/detail/merge.h
@@ -241,8 +241,8 @@ merge(execution_policy<Derived>& policy,
(__merge::merge<Derived, KeysIt1, KeysIt2, ResultIt, CompareOp>)
);
#else
-
- return __merge::merge(policy,
+ return __merge::merge(
+ policy,
keys1_first,
keys1_last,
keys2_first,
@@ -277,27 +277,6 @@ merge(execution_policy<Derived>& policy,
#else
return workaround::seq(policy, keys1_first, keys1_last, keys2_first, keys2_last, result, compare_op);
#endif
-
-
-// ResultIt ret = result;
-// THRUST_HIP_PRESERVE_KERNELS_WORKAROUND(
-// (__merge::merge<Derived, KeysIt1, KeysIt2, ResultIt, CompareOp>)
-// );
-// #if __THRUST_HAS_HIPRT__
-// ret = __merge::merge(
-// policy, keys1_first, keys1_last, keys2_first, keys2_last, result, compare_op
-// );
-// #else
-// ret = thrust::merge(cvt_to_seq(derived_cast(policy)),
-// keys1_first,
-// keys1_last,
-// keys2_first,
-// keys2_last,
-// result,
-// compare_op);
-//
-// #endif
-// return ret;
}
__thrust_exec_check_disable__ template <class Derived,
@@ -340,8 +319,8 @@ merge_by_key(execution_policy<Derived>& policy,
(__merge::merge<Derived, KeysIt1, KeysIt2, ResultIt, CompareOp>)
);
#else
-
- return __merge::merge(policy,
+ return __merge::merge(
+ policy,
keys1_first,
keys1_last,
keys2_first,

MODIFY
thrust/system/hip/detail/reduce_by_key.h
thrust/system/hip/detail/reduce_by_key.h
@@ -167,7 +167,8 @@ reduce_by_key(execution_policy<Derived>& policy,
struct workaround
{
__host__
- static pair<KeyOutputIt, ValOutputIt> par(execution_policy<Derived>& policy,
+ static pair<KeyOutputIt, ValOutputIt> par(
+ execution_policy<Derived>& policy,
KeyInputIt keys_first,
KeyInputIt keys_last,
ValInputIt values_first,
@@ -194,7 +195,8 @@ reduce_by_key(execution_policy<Derived>& policy,
#endif
}
__device__
- static pair<KeyOutputIt, ValOutputIt> seq(execution_policy<Derived>& policy,
+ static pair<KeyOutputIt, ValOutputIt> seq(
+ execution_policy<Derived>& policy,
KeyInputIt keys_first,
KeyInputIt keys_last,
ValInputIt values_first,
@@ -214,7 +216,6 @@ reduce_by_key(execution_policy<Derived>& policy,
);
}
};
-
#if __THRUST_HAS_HIPRT__
return workaround::par(policy, keys_first, keys_last, values_first, keys_output, values_output, binary_pred, binary_op);
#else

Update copy workaround and code formatting