Fixed workaround for binary search and related algorithms
MODIFY
thrust/system/hip/detail/binary_search.h
thrust/system/hip/detail/binary_search.h
@@ -229,26 +229,41 @@ lower_bound(execution_policy<Derived>& policy,
OutputIt result,
CompareOp compare_op)
{
- auto workaround_par = [=]() mutable
+ struct workaround
+ {
+ __host__
+ static OutputIt par(execution_policy<Derived>& policy,
+ HaystackIt first,
+ HaystackIt last,
+ NeedlesIt values_first,
+ NeedlesIt values_last,
+ OutputIt result,
+ CompareOp compare_op)
{
#if __HCC__ && __HIP_DEVICE_COMPILE__
- THRUST_HIP_PRESERVE_KERNELS_WORKAROUND(
+ THRUST_HIP_PRESERVE_KERNELS_WORKAROUND
(__binary_search::lower_bound<Derived,
HaystackIt,
NeedlesIt,
OutputIt,
- CompareOp>)
- );
+ CompareOp>);
return first;
#else
return __binary_search::lower_bound(
policy, first, last, values_first, values_last, result, compare_op
);
#endif
- };
- (void)workaround_par;
+ }
+
- auto workaround_seq = [=]() mutable
+ __device__
+ static OutputIt seq(execution_policy<Derived>& policy,
+ HaystackIt first,
+ HaystackIt last,
+ NeedlesIt values_first,
+ NeedlesIt values_last,
+ OutputIt result,
+ CompareOp compare_op)
{
return thrust::lower_bound(cvt_to_seq(derived_cast(policy)),
first,
@@ -257,14 +272,17 @@ lower_bound(execution_policy<Derived>& policy,
values_last,
result,
compare_op);
+ }
};
- (void)workaround_seq;
#if __THRUST_HAS_HIPRT__
- return workaround_par();
+ return workaround::par(policy, first, last, values_first, values_last, result, compare_op);
#else
- return workaround_seq();
+ return workaround::seq(policy, first, last, values_first, values_last, result, compare_op);
#endif
+
+
+
}
template <class Derived, class HaystackIt, class NeedlesIt, class OutputIt>
@@ -291,26 +309,42 @@ upper_bound(execution_policy<Derived>& policy,
OutputIt result,
CompareOp compare_op)
{
- auto workaround_par = [=]() mutable
+ struct workaround
+ {
+ __host__
+ static OutputIt par(execution_policy<Derived>& policy,
+ HaystackIt first,
+ HaystackIt last,
+ NeedlesIt values_first,
+ NeedlesIt values_last,
+ OutputIt result,
+ CompareOp compare_op)
{
#if __HCC__ && __HIP_DEVICE_COMPILE__
- THRUST_HIP_PRESERVE_KERNELS_WORKAROUND(
+ THRUST_HIP_PRESERVE_KERNELS_WORKAROUND
(__binary_search::upper_bound<Derived,
HaystackIt,
NeedlesIt,
OutputIt,
- CompareOp>)
- );
+ CompareOp>);
return first;
+
#else
return __binary_search::upper_bound(
policy, first, last, values_first, values_last, result, compare_op
);
#endif
- };
- (void)workaround_par;
+ }
- auto workaround_seq = [=]() mutable
+
+ __device__
+ static OutputIt seq(execution_policy<Derived>& policy,
+ HaystackIt first,
+ HaystackIt last,
+ NeedlesIt values_first,
+ NeedlesIt values_last,
+ OutputIt result,
+ CompareOp compare_op)
{
return thrust::upper_bound(cvt_to_seq(derived_cast(policy)),
first,
@@ -319,14 +353,15 @@ upper_bound(execution_policy<Derived>& policy,
values_last,
result,
compare_op);
+ }
};
- (void)workaround_seq;
#if __THRUST_HAS_HIPRT__
- return workaround_par();
+ return workaround::par(policy, first, last, values_first, values_last, result, compare_op);
#else
- return workaround_seq();
+ return workaround::seq(policy, first, last, values_first, values_last, result, compare_op);
#endif
+
}
template <class Derived, class HaystackIt, class NeedlesIt, class OutputIt>
@@ -353,26 +388,42 @@ binary_search(execution_policy<Derived>& policy,
OutputIt result,
CompareOp compare_op)
{
- auto workaround_par = [=]() mutable
+ struct workaround
+ {
+ __host__
+ static OutputIt par(execution_policy<Derived>& policy,
+ HaystackIt first,
+ HaystackIt last,
+ NeedlesIt values_first,
+ NeedlesIt values_last,
+ OutputIt result,
+ CompareOp compare_op)
{
#if __HCC__ && __HIP_DEVICE_COMPILE__
- THRUST_HIP_PRESERVE_KERNELS_WORKAROUND(
- (__binary_search::binary_search<Derived,
+ THRUST_HIP_PRESERVE_KERNELS_WORKAROUND
+ (__binary_search::upper_bound<Derived,
HaystackIt,
NeedlesIt,
OutputIt,
- CompareOp>)
- );
+ CompareOp>);
return first;
+
#else
return __binary_search::binary_search(
policy, first, last, values_first, values_last, result, compare_op
);
#endif
- };
- (void)workaround_par;
+ }
- auto workaround_seq = [=]() mutable
+
+ __device__
+ static OutputIt seq(execution_policy<Derived>& policy,
+ HaystackIt first,
+ HaystackIt last,
+ NeedlesIt values_first,
+ NeedlesIt values_last,
+ OutputIt result,
+ CompareOp compare_op)
{
return thrust::binary_search(cvt_to_seq(derived_cast(policy)),
first,
@@ -381,14 +432,15 @@ binary_search(execution_policy<Derived>& policy,
values_last,
result,
compare_op);
+ }
};
- (void)workaround_seq;
#if __THRUST_HAS_HIPRT__
- return workaround_par();
+ return workaround::par(policy, first, last, values_first, values_last, result, compare_op);
#else
- return workaround_seq();
+ return workaround::seq(policy, first, last, values_first, values_last, result, compare_op);
#endif
+
}
template <class Derived, class HaystackIt, class NeedlesIt, class OutputIt>
@@ -423,16 +475,22 @@ HaystackIt lower_bound(execution_policy<Derived>& policy,
using values_type = typename thrust::detail::temporary_array<T, Derived>;
using results_type = typename thrust::detail::temporary_array<difference_type, Derived>;
- auto workaround_par = [=]() mutable
+ struct workaround
+ {
+ __host__
+ static HaystackIt par(execution_policy<Derived>& policy,
+ HaystackIt first,
+ HaystackIt last,
+ const T& value,
+ CompareOp compare_op)
{
#if __HCC__ && __HIP_DEVICE_COMPILE__
- THRUST_HIP_PRESERVE_KERNELS_WORKAROUND(
+ THRUST_HIP_PRESERVE_KERNELS_WORKAROUND
(__binary_search::lower_bound<Derived,
HaystackIt,
typename values_type::iterator,
typename results_type::iterator,
- CompareOp>)
- );
+ CompareOp>);
return first;
#else
values_type values(policy, 1);
@@ -445,11 +503,17 @@ HaystackIt lower_bound(execution_policy<Derived>& policy,
);
return first + result[0];
+
#endif
- };
- (void)workaround_par;
+ }
+
- auto workaround_seq = [=]() mutable
+ __device__
+ static HaystackIt seq(execution_policy<Derived>& policy,
+ HaystackIt first,
+ HaystackIt last,
+ const T& value,
+ CompareOp compare_op)
{
difference_type result;
thrust::lower_bound(cvt_to_seq(derived_cast(policy)),
@@ -460,14 +524,15 @@ HaystackIt lower_bound(execution_policy<Derived>& policy,
&result,
compare_op);
return first + result;
+ }
};
- (void)workaround_seq;
#if __THRUST_HAS_HIPRT__
- return workaround_par();
+ return workaround::par(policy, first, last, value, compare_op);
#else
- return workaround_seq();
+ return workaround::seq(policy, first, last, value, compare_op);
#endif
+
}
template<typename Derived, typename HaystackIt, typename T, typename CompareOp>
@@ -482,17 +547,24 @@ HaystackIt upper_bound(execution_policy<Derived>& policy,
using values_type = typename thrust::detail::temporary_array<T, Derived>;
using results_type = typename thrust::detail::temporary_array<difference_type, Derived>;
- auto workaround_par = [=]() mutable
+ struct workaround
+ {
+ __host__
+ static HaystackIt par(execution_policy<Derived>& policy,
+ HaystackIt first,
+ HaystackIt last,
+ const T& value,
+ CompareOp compare_op)
{
#if __HCC__ && __HIP_DEVICE_COMPILE__
- THRUST_HIP_PRESERVE_KERNELS_WORKAROUND(
+ THRUST_HIP_PRESERVE_KERNELS_WORKAROUND
(__binary_search::upper_bound<Derived,
HaystackIt,
typename values_type::iterator,
typename results_type::iterator,
- CompareOp>)
- );
+ CompareOp>);
return first;
+
#else
values_type values(policy, 1);
results_type result(policy, 1);
@@ -504,11 +576,17 @@ HaystackIt upper_bound(execution_policy<Derived>& policy,
);
return first + result[0];
+
#endif
- };
- (void)workaround_par;
+ }
- auto workaround_seq = [=]() mutable
+
+ __device__
+ static HaystackIt seq(execution_policy<Derived>& policy,
+ HaystackIt first,
+ HaystackIt last,
+ const T& value,
+ CompareOp compare_op)
{
difference_type result;
thrust::upper_bound(cvt_to_seq(derived_cast(policy)),
@@ -519,13 +597,13 @@ HaystackIt upper_bound(execution_policy<Derived>& policy,
&result,
compare_op);
return first + result;
+ }
};
- (void)workaround_seq;
#if __THRUST_HAS_HIPRT__
- return workaround_par();
+ return workaround::par(policy, first, last, value, compare_op);
#else
- return workaround_seq();
+ return workaround::seq(policy, first, last, value, compare_op);
#endif
}
@@ -540,16 +618,22 @@ bool binary_search(execution_policy<Derived>& policy,
using values_type = typename thrust::detail::temporary_array<T, Derived>;
using results_type = typename thrust::detail::temporary_array<int, Derived>;
- auto workaround_par = [=]() mutable
+ struct workaround
+ {
+ __host__
+ static bool par(execution_policy<Derived>& policy,
+ HaystackIt first,
+ HaystackIt last,
+ const T& value,
+ CompareOp compare_op)
{
#if __HCC__ && __HIP_DEVICE_COMPILE__
- THRUST_HIP_PRESERVE_KERNELS_WORKAROUND(
- (__binary_search::binary_search<Derived,
+ THRUST_HIP_PRESERVE_KERNELS_WORKAROUND
+ (__binary_search::lower_bound<Derived,
HaystackIt,
typename values_type::iterator,
typename results_type::iterator,
- CompareOp>)
- );
+ CompareOp>);
return first;
#else
values_type values(policy, 1);
@@ -557,19 +641,24 @@ bool binary_search(execution_policy<Derived>& policy,
values[0] = value;
- __binary_search::binary_search(
+ __binary_search::lower_bound(
policy, first, last, values.begin(), values.end(), result.begin(), compare_op
);
return result[0] != 0;
#endif
- };
- (void)workaround_par;
+ }
+
- auto workaround_seq = [=]() mutable
+ __device__
+ static bool seq(execution_policy<Derived>& policy,
+ HaystackIt first,
+ HaystackIt last,
+ const T& value,
+ CompareOp compare_op)
{
bool result;
- thrust::binary_search(cvt_to_seq(derived_cast(policy)),
+ thrust::lower_bound(cvt_to_seq(derived_cast(policy)),
first,
last,
&value,
@@ -577,13 +666,13 @@ bool binary_search(execution_policy<Derived>& policy,
&result,
compare_op);
return result;
+ }
};
- (void)workaround_seq;
#if __THRUST_HAS_HIPRT__
- return workaround_par();
+ return workaround::par(policy, first, last, value, compare_op);
#else
- return workaround_seq();
+ return workaround::seq(policy, first, last, value, compare_op);
#endif
}

Fixed workaround for binary search and related algorithms