Update thrust system hip detail fill to inner_product
MODIFY
thrust/system/hip/detail/fill.h
thrust/system/hip/detail/fill.h
#include <thrust/system/hip/detail/transform.h>
BEGIN_NS_THRUST
-namespace hip_rocprim {
+namespace hip_rocprim
+{
template <class Derived, class OutputIterator, class Size, class T>
-OutputIterator __host__ __device__
+OutputIterator THRUST_HIP_FUNCTION
fill_n(execution_policy<Derived>& policy,
OutputIterator first,
Size count,
@@ -52,7 +53,7 @@ fill_n(execution_policy<Derived>& policy,
} // func fill_n
template <class Derived, class ForwardIterator, class T>
-void __host__ __device__
+void THRUST_HIP_FUNCTION
fill(execution_policy<Derived>& policy,
ForwardIterator first,
ForwardIterator last,

MODIFY
thrust/system/hip/detail/find.h
thrust/system/hip/detail/find.h
******************************************************************************/
#pragma once
-
#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
#include <thrust/system/hip/config.h>
-#include <thrust/system/hip/detail/execution_policy.h>
#include <thrust/detail/minmax.h>
#include <thrust/distance.h>
+#include <thrust/system/hip/detail/execution_policy.h>
BEGIN_NS_THRUST
-namespace hip_rocprim {
+namespace hip_rocprim
+{
// XXX forward declare to circumvent circular depedency
-template <class Derived,
- class InputIt,
- class Predicate>
-InputIt __host__ __device__
+template <class Derived, class InputIt, class Predicate>
+InputIt THRUST_HIP_FUNCTION
find_if(execution_policy<Derived>& policy,
InputIt first,
InputIt last,
Predicate predicate);
-template <class Derived,
- class InputIt,
- class Predicate>
-InputIt __host__ __device__
+template <class Derived, class InputIt, class Predicate>
+InputIt THRUST_HIP_FUNCTION
find_if_not(execution_policy<Derived>& policy,
InputIt first,
InputIt last,
Predicate predicate);
-template <class Derived,
- class InputIt,
- class T>
-InputIt __host__ __device__
-find(execution_policy<Derived> &policy,
- InputIt first,
- InputIt last,
- T const& value);
+template <class Derived, class InputIt, class T>
+InputIt THRUST_HIP_FUNCTION
+find(execution_policy<Derived>& policy, InputIt first, InputIt last, T const& value);
}; // namespace hip_rocprim
END_NS_THRUST
-#include <thrust/system/hip/detail/reduce.h>
#include <thrust/iterator/zip_iterator.h>
+#include <thrust/system/hip/detail/reduce.h>
BEGIN_NS_THRUST
-namespace hip_rocprim {
-
-namespace __find_if {
-
+namespace hip_rocprim
+{
+namespace __find_if
+{
template <typename TupleType>
struct functor
{
- THRUST_HIP_DEVICE_FUNCTION TupleType
- operator()(const TupleType& lhs, const TupleType& rhs) const
+ TupleType THRUST_HIP_DEVICE_FUNCTION
+ operator()(const TupleType& lhs,
+ const TupleType& rhs) const
{
// select the smallest index among true results
if(thrust::get<0>(lhs) && thrust::get<0>(rhs))
@@ -100,11 +92,8 @@ namespace __find_if {
};
} // namespace __find_if
-template <class Derived,
- class InputIt,
- class Size,
- class Predicate>
-InputIt __host__ __device__
+template <class Derived, class InputIt, class Size, class Predicate>
+InputIt THRUST_HIP_FUNCTION
find_if_n(execution_policy<Derived>& policy,
InputIt first,
Size num_items,
@@ -113,7 +102,8 @@ find_if_n(execution_policy<Derived>& policy,
typedef typename thrust::tuple<bool, Size> result_type;
// empty sequence
- if(num_items == 0) return first;
+ if(num_items == 0)
+ return first;
// this implementation breaks up the sequence into separate intervals
// in an attempt to early-out as soon as a value is found
@@ -121,30 +111,22 @@ find_if_n(execution_policy<Derived>& policy,
// XXX compose find_if from a look-back prefix scan algorithm
// and abort kernel when the first element is found
-
// TODO incorporate sizeof(InputType) into interval_threshold and round to multiple of 32
const Size interval_threshold = 1 << 20;
const Size interval_size = (thrust::min)(interval_threshold, num_items);
// force transform_iterator output to bool
- typedef transform_input_iterator_t<bool,
- InputIt,
- Predicate>
- XfrmIterator;
- typedef thrust::tuple<XfrmIterator,
- counting_iterator_t<Size> >
- IteratorTuple;
+ typedef transform_input_iterator_t<bool, InputIt, Predicate> XfrmIterator;
+ typedef thrust::tuple<XfrmIterator, counting_iterator_t<Size>> IteratorTuple;
typedef thrust::zip_iterator<IteratorTuple> ZipIterator;
- IteratorTuple iter_tuple =
- thrust::make_tuple(XfrmIterator(first, predicate),
- counting_iterator_t<Size>(0));
+ IteratorTuple iter_tuple
+ = thrust::make_tuple(XfrmIterator(first, predicate), counting_iterator_t<Size>(0));
ZipIterator begin = thrust::make_zip_iterator(iter_tuple);
ZipIterator end = begin + num_items;
- for (ZipIterator interval_begin = begin;
- interval_begin < end;
+ for(ZipIterator interval_begin = begin; interval_begin < end;
interval_begin += interval_size)
{
ZipIterator interval_end = interval_begin + interval_size;
@@ -170,10 +152,8 @@ find_if_n(execution_policy<Derived>& policy,
return first + num_items;
}
-template <class Derived,
- class InputIt,
- class Predicate>
-InputIt __host__ __device__
+template <class Derived, class InputIt, class Predicate>
+InputIt THRUST_HIP_FUNCTION
find_if(execution_policy<Derived>& policy,
InputIt first,
InputIt last,
@@ -182,10 +162,8 @@ find_if(execution_policy<Derived>& policy,
return hip_rocprim::find_if_n(policy, first, thrust::distance(first, last), predicate);
}
-template <class Derived,
- class InputIt,
- class Predicate>
-InputIt __host__ __device__
+template <class Derived, class InputIt, class Predicate>
+InputIt THRUST_HIP_FUNCTION
find_if_not(execution_policy<Derived>& policy,
InputIt first,
InputIt last,
@@ -194,23 +172,16 @@ find_if_not(execution_policy<Derived>& policy,
return hip_rocprim::find_if(policy, first, last, detail::not1(predicate));
}
-
-template <class Derived,
- class InputIt,
- class T>
-InputIt __host__ __device__
+template <class Derived, class InputIt, class T>
+InputIt THRUST_HIP_FUNCTION
find(execution_policy<Derived>& policy,
InputIt first,
InputIt last,
T const& value)
{
- return hip_rocprim::find_if(policy,
- first,
- last,
- detail::equal_to_value<T>(value));
+ return hip_rocprim::find_if(policy, first, last, detail::equal_to_value<T>(value));
}
-
} // namespace hip_rocprim
END_NS_THRUST
#endif

MODIFY
thrust/system/hip/detail/for_each.h
thrust/system/hip/detail/for_each.h
******************************************************************************/
#pragma once
-
#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
#include <iterator>
#include <thrust/system/hip/config.h>
-#include <thrust/system/hip/detail/util.h>
-#include <thrust/system/hip/detail/parallel_for.h>
#include <thrust/detail/function.h>
#include <thrust/distance.h>
+#include <thrust/system/hip/detail/parallel_for.h>
+#include <thrust/system/hip/detail/util.h>
BEGIN_NS_THRUST
-namespace hip_rocprim {
+namespace hip_rocprim
+{
// for_each functor
template <class Input, class UnaryOp>
@@ -50,11 +50,13 @@ namespace hip_rocprim {
THRUST_HIP_FUNCTION
for_each_f(Input input, UnaryOp op)
- : input(input), op(op) {}
+ : input(input)
+ , op(op)
+ {
+ }
template <class Size>
- THRUST_HIP_FUNCTION
- void operator()(Size idx)
+ THRUST_HIP_FUNCTION void operator()(Size idx)
{
op(raw_reference_cast(input[idx]));
}
@@ -65,39 +67,27 @@ namespace hip_rocprim {
//-------------------------
// for_each_n
- template <class Derived,
- class Input,
- class Size,
- class UnaryOp>
+template <class Derived, class Input, class Size, class UnaryOp>
Input THRUST_HIP_FUNCTION
- for_each_n(execution_policy<Derived> &policy,
- Input first,
- Size count,
- UnaryOp op)
+for_each_n(execution_policy<Derived>& policy, Input first, Size count, UnaryOp op)
{
typedef detail::wrapped_function<UnaryOp, void> wrapped_t;
wrapped_t wrapped_op(op);
- hip_rocprim::parallel_for(policy,
- for_each_f<Input, wrapped_t>(first, wrapped_op),
- count);
+ hip_rocprim::parallel_for(policy, for_each_f<Input, wrapped_t>(first, wrapped_op), count);
return first + count;
}
// for_each
- template <class Derived,
- class Input,
- class UnaryOp>
+template <class Derived, class Input, class UnaryOp>
Input THRUST_HIP_FUNCTION
- for_each(execution_policy<Derived> &policy,
- Input first,
- Input last,
- UnaryOp op)
+for_each(execution_policy<Derived>& policy, Input first, Input last, UnaryOp op)
{
typedef typename iterator_traits<Input>::difference_type size_type;
size_type count = static_cast<size_type>(thrust::distance(first, last));
return hip_rocprim::for_each_n(policy, first, count, op);
}
+
} // namespace hip_rocprim
END_NS_THRUST

MODIFY
thrust/system/hip/detail/gather.h
thrust/system/hip/detail/gather.h
******************************************************************************/
#pragma once
-
#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
-#include <thrust/system/hip/detail/transform.h>
#include <thrust/iterator/permutation_iterator.h>
+#include <thrust/system/hip/detail/transform.h>
BEGIN_NS_THRUST
-namespace hip_rocprim {
+namespace hip_rocprim
+{
-template <class Derived,
- class MapIt,
- class ItemsIt,
- class ResultIt>
-ResultIt __host__ __device__
+template <class Derived, class MapIt, class ItemsIt, class ResultIt>
+ResultIt THRUST_HIP_FUNCTION
gather(execution_policy<Derived>& policy,
MapIt map_first,
MapIt map_last,
@@ -59,7 +56,7 @@ template <class Derived,
class ItemsIt,
class ResultIt,
class Predicate>
-ResultIt __host__ __device__
+ResultIt THRUST_HIP_FUNCTION
gather_if(execution_policy<Derived>& policy,
MapIt map_first,
MapIt map_last,
@@ -77,12 +74,8 @@ gather_if(execution_policy<Derived>& policy,
predicate);
}
-template <class Derived,
- class MapIt,
- class StencilIt,
- class ItemsIt,
- class ResultIt>
-ResultIt __host__ __device__
+template <class Derived, class MapIt, class StencilIt, class ItemsIt, class ResultIt>
+ResultIt THRUST_HIP_FUNCTION
gather_if(execution_policy<Derived>& policy,
MapIt map_first,
MapIt map_last,
@@ -90,13 +83,9 @@ gather_if(execution_policy<Derived>& policy,
ItemsIt items,
ResultIt result)
{
- return hip_rocprim::gather_if(policy,
- map_first,
- map_last,
- stencil,
- items,
- result,
- identity());
+ return hip_rocprim::gather_if(
+ policy, map_first, map_last, stencil, items, result, identity()
+ );
}
} // namespace hip_rocprim

MODIFY
thrust/system/hip/detail/generate.h
thrust/system/hip/detail/generate.h
#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
-#include <thrust/distance.h>
#include <iterator>
+#include <thrust/distance.h>
#include <thrust/system/hip/config.h>
#include <thrust/system/hip/detail/for_each.h>
BEGIN_NS_THRUST
-namespace hip_rocprim {
+namespace hip_rocprim
+{
// for_each functor
template <class Generator>
@@ -44,7 +45,10 @@ struct generate_f
Generator generator;
THRUST_HIP_FUNCTION
- generate_f(Generator generator_) : generator(generator_) {}
+ generate_f(Generator generator_)
+ : generator(generator_)
+ {
+ }
template <class T>
THRUST_HIP_DEVICE_FUNCTION void operator()(T const& value)
@@ -56,21 +60,18 @@ struct generate_f
// for_each_n
template <class Derived, class OutputIt, class Size, class Generator>
-OutputIt __host__ __device__
+OutputIt THRUST_HIP_FUNCTION
generate_n(execution_policy<Derived>& policy,
OutputIt result,
Size count,
Generator generator)
{
- return hip_rocprim::for_each_n(policy,
- result,
- count,
- generate_f<Generator>(generator));
+ return hip_rocprim::for_each_n(policy, result, count, generate_f<Generator>(generator));
}
// for_each
template <class Derived, class OutputIt, class Generator>
-void __host__ __device__
+void THRUST_HIP_FUNCTION
generate(execution_policy<Derived>& policy,
OutputIt first,
OutputIt last,

MODIFY
thrust/system/hip/detail/get_value.h
thrust/system/hip/detail/get_value.h
#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
#include <thrust/detail/config.h>
#include <thrust/detail/raw_pointer_cast.h>
+#include <thrust/iterator/iterator_traits.h>
#include <thrust/system/detail/adl/assign_value.h>
#include <thrust/system/hip/config.h>
#include <thrust/system/hip/detail/cross_system.h>
-#include <thrust/iterator/iterator_traits.h>
BEGIN_NS_THRUST
-namespace hip_rocprim {
-
+namespace hip_rocprim
+{
template <typename DerivedPolicy, typename Pointer>
-inline __host__ __device__
- typename thrust::iterator_value<Pointer>::type
+typename thrust::iterator_value<Pointer>::type THRUST_HIP_FUNCTION
get_value(execution_policy<DerivedPolicy>& exec, Pointer ptr)
{
typedef typename thrust::iterator_value<Pointer>::type result_type;
// STREAMHPC WORKAROUND
#if defined(THRUST_HIP_DEVICE_CODE)
-
THRUST_UNUSED_VAR(exec);
- void (*fptr)(cross_system<thrust::host_system_tag, DerivedPolicy> &, result_type *, Pointer) = assign_value;
+ void (*fptr)(cross_system<thrust::host_system_tag, DerivedPolicy>&, result_type*, Pointer)
+ = assign_value;
(void)fptr;
return *thrust::raw_pointer_cast(ptr);
-
#else
-
// when called from host code, implement with assign_value
// note that this requires a type with default constructor
result_type result;
@@ -56,11 +53,9 @@ inline __host__ __device__
assign_value(systems, &result, ptr);
return result;
-
#endif
} // end get_value()
-
} // end hip_rocprim
END_NS_THRUST

MODIFY
thrust/system/hip/detail/inner_product.h
thrust/system/hip/detail/inner_product.h
******************************************************************************/
#pragma once
-
#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
#include <iterator>
-#include <thrust/system/hip/detail/reduce.h>
#include <thrust/detail/minmax.h>
#include <thrust/distance.h>
+#include <thrust/system/hip/detail/reduce.h>
// rocPRIM includes
#include <rocprim/rocprim.hpp>
BEGIN_NS_THRUST
-namespace hip_rocprim {
+namespace hip_rocprim
+{
template <class Derived,
class InputIt1,
@@ -47,7 +47,7 @@ template <class Derived,
class T,
class ReduceOp,
class ProductOp>
-T __host__ __device__
+T THRUST_HIP_FUNCTION
inner_product(execution_policy<Derived>& policy,
InputIt1 first1,
InputIt1 last1,
@@ -58,37 +58,25 @@ inner_product(execution_policy<Derived> &policy,
{
typedef typename iterator_traits<InputIt1>::difference_type size_type;
size_type num_items = static_cast<size_type>(thrust::distance(first1, last1));
- typedef transform_pair_of_input_iterators_t<T,
- InputIt1,
- InputIt2,
- ProductOp>
+ typedef transform_pair_of_input_iterators_t<T, InputIt1, InputIt2, ProductOp>
binop_iterator_t;
- return hip_rocprim::reduce_n(policy,
- binop_iterator_t(first1, first2, product_op),
- num_items,
- init,
- reduce_op);
+ return hip_rocprim::reduce_n(
+ policy, binop_iterator_t(first1, first2, product_op), num_items, init, reduce_op
+ );
}
-template <class Derived,
- class InputIt1,
- class InputIt2,
- class T>
-T __host__ __device__
+template <class Derived, class InputIt1, class InputIt2, class T>
+T THRUST_HIP_FUNCTION
inner_product(execution_policy<Derived>& policy,
InputIt1 first1,
InputIt1 last1,
InputIt2 first2,
T init)
{
- return hip_rocprim::inner_product(policy,
- first1,
- last1,
- first2,
- init,
- plus<T>(),
- multiplies<T>());
+ return hip_rocprim::inner_product(
+ policy, first1, last1, first2, init, plus<T>(), multiplies<T>()
+ );
}
} // namespace hip_rocprim

Update thrust system hip detail fill to inner_product