Update unique workarounds
MODIFY
thrust/system/hip/detail/unique.h
thrust/system/hip/detail/unique.h
@@ -142,7 +142,8 @@ unique_copy(execution_policy<Derived>& policy,
OutputIt result,
BinaryPred binary_pred)
{
-
+ // struct workaround is required for HIP-clang
+ // THRUST_HIP_PRESERVE_KERNELS_WORKAROUND is required for HCC
struct workaround
{
__host__
@@ -176,9 +177,6 @@ unique_copy(execution_policy<Derived>& policy,
#else
return workaround::seq(policy, first, last, result, binary_pred);
#endif
-
-
-
}
template <class Derived, class InputIt, class OutputIt>
@@ -196,16 +194,40 @@ unique(execution_policy<Derived>& policy,
InputIt last,
BinaryPred binary_pred)
{
- InputIt ret = first;
+ // struct workaround is required for HIP-clang
+ // THRUST_HIP_PRESERVE_KERNELS_WORKAROUND is required for HCC
+ struct workaround
+ {
+ __host__
+ static InputIt par(execution_policy<Derived>& policy,
+ InputIt first,
+ InputIt last,
+ BinaryPred binary_pred)
+ {
+ #if __HCC__ && __HIP_DEVICE_COMPILE__
THRUST_HIP_PRESERVE_KERNELS_WORKAROUND(
(unique_copy<Derived, InputIt, InputIt, BinaryPred>)
);
+ #else
+ return hip_rocprim::unique_copy(policy, first, last, first, binary_pred);
+ #endif
+ }
+ __device__
+ static InputIt seq(execution_policy<Derived>& policy,
+ InputIt first,
+ InputIt last,
+ BinaryPred binary_pred)
+ {
+ return thrust::unique(
+ cvt_to_seq(derived_cast(policy)), first, last, binary_pred
+ );
+ }
+ };
#if __THRUST_HAS_HIPRT__
- ret = hip_rocprim::unique_copy(policy, first, last, first, binary_pred);
+ return workaround::par(policy, first, last, binary_pred);
#else
- ret = thrust::unique(cvt_to_seq(derived_cast(policy)), first, last, binary_pred);
+ return workaround::seq(policy, first, last, binary_pred);
#endif
- return ret;
}
template <class Derived, class InputIt>

MODIFY
thrust/system/hip/detail/unique_by_key.h
thrust/system/hip/detail/unique_by_key.h
@@ -182,7 +182,6 @@ pair<KeyOutputIt, ValOutputIt>
ValOutputIt values_result,
BinaryPred binary_pred)
{
-
struct workaround
{
__host__
@@ -196,15 +195,16 @@ pair<KeyOutputIt, ValOutputIt>
{
#if __HCC__ && __HIP_DEVICE_COMPILE__
THRUST_HIP_PRESERVE_KERNELS_WORKAROUND(
- __unique_by_key::unique_by_key<Derived,
+ (__unique_by_key::unique_by_key<
+ Derived,
KeyInputIt,
ValInputIt,
KeyOutputIt,
ValOutputIt,
- BinaryPred>));
+ BinaryPred>)
+ );
#else
- return __unique_by_key::unique_by_key
- (
+ return __unique_by_key::unique_by_key(
policy,
keys_first,
keys_last,
@@ -224,8 +224,7 @@ pair<KeyOutputIt, ValOutputIt>
ValOutputIt values_result,
BinaryPred binary_pred)
{
- return thrust::unique_by_key_copy
- (
+ return thrust::unique_by_key_copy(
cvt_to_seq(derived_cast(policy)),
keys_first,
keys_last,
@@ -236,14 +235,11 @@ pair<KeyOutputIt, ValOutputIt>
);
}
};
-
#if __THRUST_HAS_HIPRT__
return workaround::par(policy, keys_first, keys_last, values_first, keys_result, values_result, binary_pred);
#else
return workaround::seq(policy, keys_first, keys_last, values_first, keys_result, values_result, binary_pred);
#endif
-
-
}
template <class Derived,
@@ -277,7 +273,17 @@ unique_by_key(execution_policy<Derived>& policy,
ValInputIt values_first,
BinaryPred binary_pred)
{
- pair<KeyInputIt, ValInputIt> ret = thrust::make_pair(keys_first, values_first);
+ struct workaround
+ {
+ __host__
+ static pair<KeyInputIt, ValInputIt>par(
+ execution_policy<Derived>& policy,
+ KeyInputIt keys_first,
+ KeyInputIt keys_last,
+ ValInputIt values_first,
+ BinaryPred binary_pred)
+ {
+ #if __HCC__ && __HIP_DEVICE_COMPILE__
THRUST_HIP_PRESERVE_KERNELS_WORKAROUND(
(hip_rocprim::unique_by_key_copy<Derived,
KeyInputIt,
@@ -285,16 +291,36 @@ unique_by_key(execution_policy<Derived>& policy,
KeyInputIt,
ValInputIt>)
);
-#if __THRUST_HAS_HIPRT__
- ret = hip_rocprim::unique_by_key_copy(
- policy, keys_first, keys_last, values_first, keys_first, values_first, binary_pred
- );
#else
- ret = thrust::unique_by_key(
+ return hip_rocprim::unique_by_key_copy(
+ policy,
+ keys_first,
+ keys_last,
+ values_first,
+ keys_first,
+ values_first,
+ binary_pred
+ );
+ #endif
+ }
+ __device__
+ static pair<KeyInputIt, ValInputIt> seq(
+ execution_policy<Derived>& policy,
+ KeyInputIt keys_first,
+ KeyInputIt keys_last,
+ ValInputIt values_first,
+ BinaryPred binary_pred)
+ {
+ return thrust::unique_by_key(
cvt_to_seq(derived_cast(policy)), keys_first, keys_last, values_first, binary_pred
);
+ }
+ };
+ #if __THRUST_HAS_HIPRT__
+ return workaround::par(policy, keys_first, keys_last, values_first, binary_pred);
+ #else
+ return workaround::seq(policy, keys_first, keys_last, values_first, binary_pred);
#endif
- return ret;
}
template <class Derived, class KeyInputIt, class ValInputIt>

Update unique workarounds