 #include <thrust/execution_policy.h>
 
 
-template<typename Iterator1, typename Iterator2>
+template<typename ExecutionPolicy, typename Iterator1, typename Iterator2>
 __global__
-void unique_kernel(Iterator1 first, Iterator1 last, Iterator2 result)
+void unique_kernel(ExecutionPolicy exec, Iterator1 first, Iterator1 last, Iterator2 result)
 {
-  *result = thrust::unique(thrust::seq, first, last);
+  *result = thrust::unique(exec, first, last);
 }
 
 
-template<typename Iterator1, typename BinaryPredicate, typename Iterator2>
+template<typename ExecutionPolicy, typename Iterator1, typename BinaryPredicate, typename Iterator2>
 __global__
-void unique_kernel(Iterator1 first, Iterator1 last, BinaryPredicate pred, Iterator2 result)
+void unique_kernel(ExecutionPolicy exec, Iterator1 first, Iterator1 last, BinaryPredicate pred, Iterator2 result)
 {
-  *result = thrust::unique(thrust::seq, first, last, pred);
+  *result = thrust::unique(exec, first, last, pred);
 }
 
 

  ASSERT_EQUAL(data[1], 20);
  ASSERT_EQUAL(data[2], 31);
}
DECLARE_UNITTEST(TestUniqueDeviceSeq);


template<typename Iterator1, typename Iterator2, typename Iterator3>
__global__
void unique_copy_kernel(Iterator1 first, Iterator1 last, Iterator2 result1, Iterator3 result2)
{
  *result2 = thrust::unique_copy(thrust::seq, first, last, result1);
}


template<typename Iterator1, typename Iterator2, typename BinaryPredicate, typename Iterator3>
__global__
void unique_copy_kernel(Iterator1 first, Iterator1 last, Iterator2 result1, BinaryPredicate pred, Iterator3 result2)
{
  *result2 = thrust::unique_copy(thrust::seq, first, last, result1, pred);
}


void TestUniqueCopyDeviceSeq()
{
  typedef thrust::device_vector<int> Vector;
  typedef Vector::value_type T;

  thrust::device_vector<Vector::iterator> new_last_vec(1);
  Vector::iterator new_last;
  
  unique_kernel<<<1,1>>>(exec, data.begin(), data.end(), new_last_vec.begin());
  new_last = new_last_vec[0];

  ASSERT_EQUAL(new_last - data.begin(), 7);
