 };
 
 
-template<typename T>
-void TestPartitionPointDeviceSeq(size_t n)
+template<typename T, typename ExecutionPolicy>
+void TestPartitionPointDevice(ExecutionPolicy exec, size_t n)
 {
   thrust::device_vector<T> v = unittest::random_integers<T>(n);
   typedef typename thrust::device_vector<T>::iterator iterator;

  iterator ref = thrust::stable_partition(v.begin(), v.end(), is_even<T>());

  thrust::device_vector<iterator> result(1);
  partition_point_kernel<<<1,1>>>(v.begin(), v.end(), is_even<T>(), result.begin());

  ASSERT_EQUAL(ref - v.begin(), (iterator)result[0] - v.begin());
}
DECLARE_VARIABLE_UNITTEST(TestPartitionPointDeviceSeq);


#include <thrust/execution_policy.h>


template<typename ExecutionPolicy, typename Iterator1, typename Predicate, typename Iterator2>
__global__
void partition_point_kernel(ExecutionPolicy exec, Iterator1 first, Iterator1 last, Predicate pred, Iterator2 result)
{
  *result = thrust::partition_point(exec, first, last, pred);
}


