Added hip::par workaround for merge, partition, reduce, and adjacent_difference
MODIFY
test/test_adjacent_difference.cpp
test/test_adjacent_difference.cpp
@@ -247,7 +247,7 @@ void AdjacentDifferenceKernel(int const N, int* array)
thrust::device_ptr<int> begin(array);
thrust::device_ptr<int> end(array + N);
// thrust::advance(begin,2);
- thrust::adjacent_difference(thrust::seq, begin,end,begin);
+ thrust::adjacent_difference(thrust::hip::par, begin,end,begin);
}
}

MODIFY
test/test_binary_search.cpp
test/test_binary_search.cpp
@@ -479,7 +479,7 @@ void BinarySearchKernel(int const N, int* in_array, int*result_array, int search
{
thrust::device_ptr<int> begin(in_array);
thrust::device_ptr<int> end(in_array + N);
- result_array[search_value]=thrust::binary_search(thrust::seq, begin,end,search_value);
+ result_array[search_value]=thrust::binary_search(thrust::hip::par, begin,end,search_value);
}
}

MODIFY
test/test_find.cpp
test/test_find.cpp
@@ -395,7 +395,7 @@ void FindKernel(int const N, int* in_array, int value, int *out_array)
thrust::device_ptr<int> in_begin(in_array);
thrust::device_ptr<int> in_end(in_array + N);
- auto x = thrust::find(thrust::seq, in_begin, in_end,value);
+ auto x = thrust::find(thrust::hip::par, in_begin, in_end,value);
out_array[0] = x - in_begin;
}
}

MODIFY
test/test_for_each.cpp
test/test_for_each.cpp
@@ -96,7 +96,7 @@ THRUST_HIP_LAUNCH_BOUNDS_DEFAULT
void simple_test_kernel(F func, int size)
{
// (void) func; (void) size;
- thrust::for_each(thrust::seq,
+ thrust::for_each(thrust::hip::par,
thrust::make_counting_iterator<int>(0),
thrust::make_counting_iterator<int>(size),
func);

MODIFY
test/test_merge.cpp
test/test_merge.cpp
@@ -290,7 +290,7 @@ void MergeKernel(int const N, T* inA_array, T* inB_array, T *out_array)
thrust::device_ptr<int> inB_end(inB_array + N);
thrust::device_ptr<int> out_begin(out_array);
- thrust::merge(thrust::seq, inA_begin, inA_end, inB_begin, inB_end,out_begin);
+ thrust::merge(thrust::hip::par, inA_begin, inA_end, inB_begin, inB_end,out_begin);
}
}

MODIFY
test/test_merge_by_key.cpp
test/test_merge_by_key.cpp
@@ -460,7 +460,7 @@ void MergeByKeyKernel(int const N, T* keys_A, T* keys_B, T * values_A, T* values
thrust::device_ptr<int> values_result_begin(values_result);
- thrust::merge_by_key(thrust::seq, keyA_begin, keyA_end,
+ thrust::merge_by_key(thrust::hip::par, keyA_begin, keyA_end,
keyB_begin, keyB_end,
valuesA_begin, valuesB_begin,
keys_result_begin,

MODIFY
test/test_partition.cpp
test/test_partition.cpp
@@ -1805,7 +1805,7 @@ void PartitionKernel(int const N, int* array)
{
thrust::device_ptr<int> begin(array);
thrust::device_ptr<int> end(array + N);
- thrust::partition(thrust::seq, begin, end,is_even<int>());
+ thrust::partition(thrust::hip::par, begin, end,is_even<int>());
}
}

MODIFY
test/test_reduce.cpp
test/test_reduce.cpp
@@ -266,7 +266,7 @@ void ReduceKernel(int const N, int *in_array, int *result)
thrust::device_ptr<int> in_begin(in_array);
thrust::device_ptr<int> in_end(in_array + N);
- result[0]= thrust::reduce(thrust::seq, in_begin, in_end);
+ result[0]= thrust::reduce(thrust::hip::par, in_begin, in_end);
}
}

MODIFY
thrust/system/hip/detail/adjacent_difference.h
thrust/system/hip/detail/adjacent_difference.h
@@ -414,18 +414,54 @@ adjacent_difference(execution_policy<Derived>& policy,
OutputIt result,
BinaryOp binary_op)
{
- OutputIt ret = result;
+ // struct workaround is required for HIP-clang
+ // THRUST_HIP_PRESERVE_KERNELS_WORKAROUND is required for HCC
+ struct workaround
+ {
+ __host__
+ static void par(execution_policy<Derived>& policy,
+ InputIt first,
+ InputIt last,
+ OutputIt& result,
+ BinaryOp binary_op)
+ {
+ #if __HCC__ && __HIP_DEVICE_COMPILE__
THRUST_HIP_PRESERVE_KERNELS_WORKAROUND(
(__adjacent_difference::adjacent_difference<Derived, InputIt, OutputIt, BinaryOp>)
);
+ #else
+ result = __adjacent_difference::adjacent_difference(
+ policy,
+ first,
+ last,
+ result,
+ binary_op
+ );
+ #endif
+ }
+ __device__
+ static void seq(execution_policy<Derived>& policy,
+ InputIt first,
+ InputIt last,
+ OutputIt& result,
+ BinaryOp binary_op)
+ {
+ result = thrust::adjacent_difference(
+ cvt_to_seq(derived_cast(policy)),
+ first,
+ last,
+ result,
+ binary_op
+ );
+ }
+ };
#if __THRUST_HAS_HIPRT__
- ret = __adjacent_difference::adjacent_difference(policy, first, last, result, binary_op);
- #else // __THRUST_HAS_HIPRT__
- ret = thrust::adjacent_difference(
- cvt_to_seq(derived_cast(policy)), first, last, result, binary_op);
- #endif // __THRUST_HAS_HIPRT__
+ workaround::par(policy, first, last, result, binary_op);
+ #else
+ workaround::seq(policy, first, last, result, binary_op);
+ #endif
- return ret;
+ return result;
}
template <class Derived, class InputIt, class OutputIt>

MODIFY
thrust/system/hip/detail/merge.h
thrust/system/hip/detail/merge.h
@@ -224,25 +224,80 @@ merge(execution_policy<Derived>& policy,
CompareOp compare_op)
{
- ResultIt ret = result;
+
+ struct workaround
+ {
+ __host__
+ static ResultIt par(execution_policy<Derived>& policy,
+ KeysIt1 keys1_first,
+ KeysIt1 keys1_last,
+ KeysIt2 keys2_first,
+ KeysIt2 keys2_last,
+ ResultIt result,
+ CompareOp compare_op)
+ {
+ #if __HCC__ && __HIP_DEVICE_COMPILE__
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)),
+
+ return __merge::merge(policy,
keys1_first,
keys1_last,
keys2_first,
keys2_last,
result,
- compare_op);
-
+ compare_op
+ );
+ #endif
+ }
+ __device__
+ static ResultIt seq(execution_policy<Derived>& policy,
+ KeysIt1 keys1_first,
+ KeysIt1 keys1_last,
+ KeysIt2 keys2_first,
+ KeysIt2 keys2_last,
+ ResultIt result,
+ CompareOp compare_op)
+ {
+ return thrust::merge(
+ cvt_to_seq(derived_cast(policy)),
+ keys1_first,
+ keys1_last,
+ keys2_first,
+ keys2_last,
+ result,
+ compare_op
+ );
+ }
+ };
+ #if __THRUST_HAS_HIPRT__
+ return workaround::par(policy, keys1_first, keys1_last, keys2_first, keys2_last, result, compare_op);
+ #else
+ return workaround::seq(policy, keys1_first, keys1_last, keys2_first, keys2_last, result, compare_op);
#endif
- return ret;
+
+
+// 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,
@@ -265,18 +320,27 @@ merge_by_key(execution_policy<Derived>& policy,
ItemsOutputIt items_result,
CompareOp compare_op)
{
- pair<KeysOutputIt, ItemsOutputIt> ret = thrust::make_pair(keys_result, items_result);
+
+ struct workaround
+ {
+ __host__
+ static pair<KeysOutputIt, ItemsOutputIt> par(execution_policy<Derived>& policy,
+ KeysIt1 keys1_first,
+ KeysIt1 keys1_last,
+ KeysIt2 keys2_first,
+ KeysIt2 keys2_last,
+ ItemsIt1 items1_first,
+ ItemsIt2 items2_first,
+ KeysOutputIt keys_result,
+ ItemsOutputIt items_result,
+ CompareOp compare_op)
+ {
+ #if __HCC__ && __HIP_DEVICE_COMPILE__
THRUST_HIP_PRESERVE_KERNELS_WORKAROUND(
- (__merge::merge<Derived,
- KeysIt1,
- KeysIt2,
- ItemsIt1,
- ItemsIt2,
- KeysOutputIt,
- ItemsOutputIt,
- CompareOp>)
+ (__merge::merge<Derived, KeysIt1, KeysIt2, ResultIt, CompareOp>)
);
-#if __THRUST_HAS_HIPRT__
+ #else
+
return __merge::merge(policy,
keys1_first,
keys1_last,
@@ -286,10 +350,24 @@ merge_by_key(execution_policy<Derived>& policy,
items2_first,
keys_result,
items_result,
- compare_op);
-
-#else
- ret = thrust::merge_by_key(cvt_to_seq(derived_cast(policy)),
+ compare_op
+ );
+ #endif
+ }
+ __device__
+ static pair<KeysOutputIt, ItemsOutputIt> seq(execution_policy<Derived>& policy,
+ KeysIt1 keys1_first,
+ KeysIt1 keys1_last,
+ KeysIt2 keys2_first,
+ KeysIt2 keys2_last,
+ ItemsIt1 items1_first,
+ ItemsIt2 items2_first,
+ KeysOutputIt keys_result,
+ ItemsOutputIt items_result,
+ CompareOp compare_op)
+ {
+ return thrust::merge_by_key(
+ cvt_to_seq(derived_cast(policy)),
keys1_first,
keys1_last,
keys2_first,
@@ -298,9 +376,53 @@ merge_by_key(execution_policy<Derived>& policy,
items2_first,
keys_result,
items_result,
- compare_op);
+ compare_op
+ );
+ }
+ };
+ #if __THRUST_HAS_HIPRT__
+ return workaround::par(policy, keys1_first, keys1_last, keys2_first, keys2_last, items1_first, items2_first, keys_result, items_result, compare_op);
+ #else
+ return workaround::seq(policy, keys1_first, keys1_last, keys2_first, keys2_last, items1_first, items2_first, keys_result, items_result, compare_op);
#endif
- return ret;
+
+
+// pair<KeysOutputIt, ItemsOutputIt> ret = thrust::make_pair(keys_result, items_result);
+// THRUST_HIP_PRESERVE_KERNELS_WORKAROUND(
+// (__merge::merge<Derived,
+// KeysIt1,
+// KeysIt2,
+// ItemsIt1,
+// ItemsIt2,
+// KeysOutputIt,
+// ItemsOutputIt,
+// CompareOp>)
+// );
+// #if __THRUST_HAS_HIPRT__
+// return __merge::merge(policy,
+// keys1_first,
+// keys1_last,
+// keys2_first,
+// keys2_last,
+// items1_first,
+// items2_first,
+// keys_result,
+// items_result,
+// compare_op);
+//
+// #else
+// ret = thrust::merge_by_key(cvt_to_seq(derived_cast(policy)),
+// keys1_first,
+// keys1_last,
+// keys2_first,
+// keys2_last,
+// items1_first,
+// items2_first,
+// keys_result,
+// items_result,
+// compare_op);
+// #endif
+// return ret;
}
__thrust_exec_check_disable__ template <class Derived,

MODIFY
thrust/system/hip/detail/partition.h
thrust/system/hip/detail/partition.h
@@ -556,16 +556,59 @@ partition(execution_policy<Derived>& policy,
Iterator last,
Predicate predicate)
{
- Iterator ret = first;
+
+ struct workaround
+ {
+ __host__
+ static Iterator par(execution_policy<Derived>& policy,
+ Iterator first,
+ Iterator last,
+ Predicate predicate)
+ {
+ #if __HCC__ && __HIP_DEVICE_COMPILE__
THRUST_HIP_PRESERVE_KERNELS_WORKAROUND(
- (__partition::partition_inplace<Derived, Iterator, Predicate>)
+ (__partition::partition<Derived, Iterator, Iterator, StencilIt, Predicate>)
);
+ #else
+
+ return __partition::partition_inplace(policy,
+ first,
+ last,
+ predicate
+ );
+ #endif
+ }
+ __device__
+ static Iterator seq(execution_policy<Derived>& policy,
+ Iterator first,
+ Iterator last,
+ Predicate predicate)
+ {
+ return thrust::partition(
+ cvt_to_seq(derived_cast(policy)),
+ first,
+ last,
+ predicate
+ );
+ }
+ };
#if __THRUST_HAS_HIPRT__
- ret = __partition::partition_inplace(policy, first, last, predicate);
-#else // __THRUST_HAS_HIPRT__
- ret = thrust::partition(cvt_to_seq(derived_cast(policy)), first, last, predicate);
-#endif // __THRUST_HAS_HIPRT__
- return ret;
+ return workaround::par(policy, first, last, predicate);
+ #else
+ return workaround::seq(policy, first, last, predicate);
+ #endif
+//
+//
+// Iterator ret = first;
+// THRUST_HIP_PRESERVE_KERNELS_WORKAROUND(
+// (__partition::partition_inplace<Derived, Iterator, Predicate>)
+// );
+// #if __THRUST_HAS_HIPRT__
+// ret = __partition::partition_inplace(policy, first, last, predicate);
+// #else // __THRUST_HAS_HIPRT__
+// ret = thrust::partition(cvt_to_seq(derived_cast(policy)), first, last, predicate);
+// #endif // __THRUST_HAS_HIPRT__
+// return ret;
}
__thrust_exec_check_disable__ template <class Derived,

MODIFY
thrust/system/hip/detail/reduce.h
thrust/system/hip/detail/reduce.h
@@ -133,16 +133,63 @@ T reduce_n(execution_policy<Derived>& policy,
T init,
BinaryOp binary_op)
{
+
+ struct workaround
+ {
+ __host__
+ static T par(execution_policy<Derived>& policy,
+ InputIt first,
+ Size num_items,
+ T init,
+ BinaryOp binary_op)
+ {
+ #if __HCC__ && __HIP_DEVICE_COMPILE__
THRUST_HIP_PRESERVE_KERNELS_WORKAROUND(
(__reduce::reduce<Derived, InputIt, Size, T, BinaryOp>)
);
-#if __THRUST_HAS_HIPRT__
- return __reduce::reduce(policy, first, num_items, init, binary_op);
-#else // __THRUST_HAS_HIPRT__
+ #else
+ return __reduce::reduce(
+ policy,
+ first,
+ num_items,
+ init,
+ binary_op
+ );
+ #endif
+ }
+ __device__
+ static T seq(execution_policy<Derived>& policy,
+ InputIt first,
+ Size num_items,
+ T init,
+ BinaryOp binary_op)
+ {
return thrust::reduce(
- cvt_to_seq(derived_cast(policy)), first, first + num_items, init, binary_op
+ cvt_to_seq(derived_cast(policy)),
+ first,
+ first + num_items,
+ init,
+ binary_op
);
-#endif // __THRUST_HAS_HIPRT__
+ }
+ };
+ #if __THRUST_HAS_HIPRT__
+ return workaround::par(policy, first, num_items, init, binary_op);
+ #else
+ return workaround::seq(policy, first, num_items, init, binary_op);
+ #endif
+
+
+// THRUST_HIP_PRESERVE_KERNELS_WORKAROUND(
+// (__reduce::reduce<Derived, InputIt, Size, T, BinaryOp>)
+// );
+// #if __THRUST_HAS_HIPRT__
+// return __reduce::reduce(policy, first, num_items, init, binary_op);
+// #else // __THRUST_HAS_HIPRT__
+// return thrust::reduce(
+// cvt_to_seq(derived_cast(policy)), first, first + num_items, init, binary_op
+// );
+// #endif // __THRUST_HAS_HIPRT__
}
template <class Derived, class InputIt, class T, class BinaryOp>

Added hip::par workaround for merge, partition, reduce, and adjacent_difference