Update copy_if workaround
MODIFY
thrust/system/hip/detail/copy_if.h
thrust/system/hip/detail/copy_if.h
@@ -263,6 +263,21 @@ copy_if(execution_policy<Derived>& policy,
OutputIterator result,
Predicate pred)
{
+
+ // struct workaround is required for HIP-clang
+ // THRUST_HIP_PRESERVE_KERNELS_WORKAROUND is required for HCC
+ struct workaround
+ {
+ __host__
+ static OutputIterator par(
+ execution_policy<Derived>& policy,
+ InputIterator first,
+ InputIterator last,
+ StencilIterator stencil,
+ OutputIterator result,
+ Predicate pred)
+ {
+ #if __HCC__ && __HIP_DEVICE_COMPILE__
THRUST_HIP_PRESERVE_KERNELS_WORKAROUND(
(__copy_if::copy_if<Derived,
InputIterator,
@@ -270,11 +285,41 @@ copy_if(execution_policy<Derived>& policy,
OutputIterator,
Predicate>)
);
-#if __THRUST_HAS_HIPRT__
- return __copy_if::copy_if(policy, first, last, stencil, result, pred);
#else
+ return __copy_if::copy_if(
+ policy,
+ first,
+ last,
+ stencil,
+ result,
+ pred
+ );
+ #endif
+ }
+ __device__
+ static OutputIterator seq(
+ execution_policy<Derived>& policy,
+ InputIterator first,
+ InputIterator last,
+ StencilIterator stencil,
+ OutputIterator result,
+ Predicate pred)
+ {
return thrust::copy_if(
- cvt_to_seq(derived_cast(policy)), first, last, stencil, result, pred);
+ cvt_to_seq(derived_cast(policy)),
+ first,
+ last,
+ stencil,
+ result,
+ pred
+ );
+ }
+ };
+
+#if __THRUST_HAS_HIPRT__
+ return workaround::par(policy, first, last, stencil, result, pred);
+#else
+ return workaround::seq(policy, first, last, stencil, result, pred);
#endif
} // func copy_if

Update copy_if workaround