   ASSERT_EQUAL(output[5], 31);
   ASSERT_EQUAL(output[6], 37);
 
-  unique_copy_kernel<<<1,1>>>(output.begin(), new_last, data.begin(), is_equal_div_10_unique<T>(), new_last_vec.begin());
+  unique_copy_kernel<<<1,1>>>(exec, output.begin(), new_last, data.begin(), is_equal_div_10_unique<T>(), new_last_vec.begin());
   new_last = new_last_vec[0];
 
   ASSERT_EQUAL(new_last - data.begin(), 3);

  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);
