Implement remove, remove_if
MODIFY
testing/CMakeLists.txt
testing/CMakeLists.txt
@@ -123,7 +123,7 @@ add_thrust_test("pair_sort")
# add_thrust_test("reduce_by_key")
# add_thrust_test("reduce")
# add_thrust_test("reduce_large")
-# add_thrust_test("remove")
+add_thrust_test("remove")
# add_thrust_test("replace")
# add_thrust_test("reverse")
# add_thrust_test("reverse_iterator")

MODIFY
testing/remove.cu
testing/remove.cu
@@ -207,10 +207,11 @@ DECLARE_VECTOR_UNITTEST(TestRemoveIfSimple);
template<typename ForwardIterator,
typename Predicate>
+__host__ __device__
ForwardIterator remove_if(my_system &system,
ForwardIterator first,
ForwardIterator,
- Predicate pred)
+ Predicate)
{
system.validate_dispatch();
return first;
@@ -230,10 +231,11 @@ DECLARE_UNITTEST(TestRemoveIfDispatchExplicit);
template<typename ForwardIterator,
typename Predicate>
+__host__ __device__
ForwardIterator remove_if(my_tag,
ForwardIterator first,
ForwardIterator,
- Predicate pred)
+ Predicate)
{
*first = 13;
return first;
@@ -288,11 +290,12 @@ DECLARE_VECTOR_UNITTEST(TestRemoveIfStencilSimple);
template<typename ForwardIterator,
typename InputIterator,
typename Predicate>
+__host__ __device__
ForwardIterator remove_if(my_system &system,
ForwardIterator first,
ForwardIterator,
InputIterator,
- Predicate pred)
+ Predicate)
{
system.validate_dispatch();
return first;
@@ -317,11 +320,12 @@ DECLARE_UNITTEST(TestRemoveIfStencilDispatchExplicit);
template<typename ForwardIterator,
typename InputIterator,
typename Predicate>
+__host__ __device__
ForwardIterator remove_if(my_tag,
ForwardIterator first,
ForwardIterator,
InputIterator,
- Predicate pred)
+ Predicate)
{
*first = 13;
return first;
@@ -372,6 +376,7 @@ DECLARE_VECTOR_UNITTEST(TestRemoveCopyIfSimple);
template<typename InputIterator,
typename OutputIterator,
typename Predicate>
+__host__ __device__
InputIterator remove_copy_if(my_system &system,
InputIterator first,
InputIterator,
@@ -401,6 +406,7 @@ DECLARE_UNITTEST(TestRemoveCopyIfDispatchExplicit);
template<typename InputIterator,
typename OutputIterator,
typename Predicate>
+__host__ __device__
InputIterator remove_copy_if(my_tag,
InputIterator first,
InputIterator,
@@ -465,6 +471,7 @@ template<typename InputIterator1,
typename InputIterator2,
typename OutputIterator,
typename Predicate>
+__host__ __device__
OutputIterator remove_copy_if(my_system &system,
InputIterator1,
InputIterator1,
@@ -497,6 +504,7 @@ template<typename InputIterator1,
typename InputIterator2,
typename OutputIterator,
typename Predicate>
+__host__ __device__
OutputIterator remove_copy_if(my_tag,
InputIterator1,
InputIterator1,
@@ -627,6 +635,8 @@ void TestRemoveCopyToDiscardIterator(const size_t n)
}
DECLARE_VARIABLE_UNITTEST(TestRemoveCopyToDiscardIterator);
+// STREAMHPC TODO fix rocPRIM select/partition for zip with discard_iterator
+#if THRUST_DEVICE_SYSTEM != THRUST_DEVICE_SYSTEM_HIP
template<typename T>
void TestRemoveCopyToDiscardIteratorZipped(const size_t n)
@@ -666,6 +676,8 @@ void TestRemoveCopyToDiscardIteratorZipped(const size_t n)
}
DECLARE_VARIABLE_UNITTEST(TestRemoveCopyToDiscardIteratorZipped);
+#endif
+
template<typename T>
void TestRemoveCopyIf(const size_t n)
@@ -759,4 +771,3 @@ void TestRemoveCopyIfStencilToDiscardIterator(const size_t n)
ASSERT_EQUAL_QUIET(reference, d_result);
}
DECLARE_VARIABLE_UNITTEST(TestRemoveCopyIfStencilToDiscardIterator);
-

MODIFY
thrust/system/detail/adl/remove.h
thrust/system/detail/adl/remove.h
#if 0
#include <thrust/system/cpp/detail/remove.h>
#include <thrust/system/cuda/detail/remove.h>
+#include <thrust/system/hip/detail/remove.h>
#include <thrust/system/omp/detail/remove.h>
#include <thrust/system/tbb/detail/remove.h>
#endif
#define __THRUST_DEVICE_SYSTEM_REMOVE_HEADER <__THRUST_DEVICE_SYSTEM_ROOT/detail/remove.h>
#include __THRUST_DEVICE_SYSTEM_REMOVE_HEADER
#undef __THRUST_DEVICE_SYSTEM_REMOVE_HEADER
-

ADD
thrust/system/hip/detail/count.h
+// STREAMHPC TODO implement

ADD
thrust/system/hip/detail/remove.h
+/******************************************************************************
+ * Copyright (c) 2016, NVIDIA CORPORATION. All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ * * Redistributions of source code must retain the above copyright
+ * notice, this list of conditions and the following disclaimer.
+ * * Redistributions in binary form must reproduce the above copyright
+ * notice, this list of conditions and the following disclaimer in the
+ * documentation and/or other materials provided with the distribution.
+ * * Neither the name of the NVIDIA CORPORATION nor the
+ * names of its contributors may be used to endorse or promote products
+ * derived from this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+ * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+ * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+ * ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
+ * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
+ * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
+ * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
+ * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ *
+ ******************************************************************************/
+#pragma once
+
+
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+#include <thrust/system/hip/detail/copy_if.h>
+
+BEGIN_NS_THRUST
+namespace hip_rocprim {
+
+// in-place
+
+// STREAMHPC TODO check if loopback implementation of select in rocPRIM really can work
+// in-place (input == output)
+
+template <class Derived,
+ class InputIt,
+ class StencilIt,
+ class Predicate>
+InputIt __host__ __device__
+remove_if(execution_policy<Derived> &policy,
+ InputIt first,
+ InputIt last,
+ StencilIt stencil,
+ Predicate predicate)
+{
+ return hip_rocprim::copy_if(policy, first, last, stencil, first, detail::not1(predicate));
+}
+
+template <class Derived,
+ class InputIt,
+ class Predicate>
+InputIt __host__ __device__
+remove_if(execution_policy<Derived> &policy,
+ InputIt first,
+ InputIt last,
+ Predicate predicate)
+{
+ return hip_rocprim::copy_if(policy, first, last, first, detail::not1(predicate));
+}
+
+template <class Derived,
+ class InputIt,
+ class T>
+InputIt __host__ __device__
+remove(execution_policy<Derived> &policy,
+ InputIt first,
+ InputIt last,
+ const T & value)
+{
+ detail::equal_to_value<T> pred(value);
+ return hip_rocprim::remove_if(policy, first, last, pred);
+}
+
+// copy
+
+template <class Derived,
+ class InputIt,
+ class StencilIt,
+ class OutputIt,
+ class Predicate>
+OutputIt __host__ __device__
+remove_copy_if(execution_policy<Derived> &policy,
+ InputIt first,
+ InputIt last,
+ StencilIt stencil,
+ OutputIt result,
+ Predicate predicate)
+{
+ return hip_rocprim::copy_if(policy, first, last, stencil, result, detail::not1(predicate));
+}
+
+template <class Derived,
+ class InputIt,
+ class OutputIt,
+ class Predicate>
+OutputIt __host__ __device__
+remove_copy_if(execution_policy<Derived> &policy,
+ InputIt first,
+ InputIt last,
+ OutputIt result,
+ Predicate predicate)
+{
+ return hip_rocprim::copy_if(policy, first, last, result, detail::not1(predicate));
+}
+
+template <class Derived,
+ class InputIt,
+ class OutputIt,
+ class T>
+OutputIt __host__ __device__
+remove_copy(execution_policy<Derived> &policy,
+ InputIt first,
+ InputIt last,
+ OutputIt result,
+ const T & value)
+{
+ detail::equal_to_value<T> pred(value);
+ return hip_rocprim::remove_copy_if(policy, first, last, result, pred);
+}
+
+} // namespace hip_rocprim
+END_NS_THRUST
+#endif

ADD
thrust/system/hip/detail/transform_reduce.h
+// STREAMHPC TODO implement

MODIFY
thrust/system/hip/execution_policy.h
thrust/system/hip/execution_policy.h
// #include <thrust/system/hip/detail/mismatch.h>
// #include <thrust/system/hip/detail/partition.h>
// #include <thrust/system/hip/detail/reduce_by_key.h>
-// #include <thrust/system/hip/detail/remove.h>
+#include <thrust/system/hip/detail/remove.h>
#include <thrust/system/hip/detail/replace.h>
// #include <thrust/system/hip/detail/reverse.h>
// #include <thrust/system/hip/detail/scatter.h>

Implement remove, remove_if