   
   ASSERT_EQUAL(h_result, d_result);
 }
+
+
+template<typename T>
+void TestRemoveCopyDeviceSeq(const size_t n)
+{
+  TestRemoveCopyDevice<T>(thrust::seq, n);
+}
 DECLARE_VARIABLE_UNITTEST(TestRemoveCopyDeviceSeq);
 
 
 template<typename T>
-void TestRemoveCopyIfDeviceSeq(const size_t n)
+void TestRemoveCopyDeviceDevice(const size_t n)
+{
+  TestRemoveCopyDevice<T>(thrust::device, n);
+}
+DECLARE_VARIABLE_UNITTEST(TestRemoveCopyDeviceDevice);
+
+
+template<typename T, typename ExecutionPolicy>
+void TestRemoveCopyIfDevice(ExecutionPolicy exec, const size_t n)
 {
   thrust::host_vector<T>   h_data = unittest::random_samples<T>(n);
   thrust::device_vector<T> d_data = h_data;

  thrust::device_vector<iterator> d_result(1);
  
  size_t h_size = thrust::remove_if(h_data.begin(), h_data.end(), is_true<T>()) - h_data.begin();
  remove_if_kernel<<<1,1>>>(d_data.begin(), d_data.end(), is_true<T>(), d_result.begin());
  size_t d_size = (iterator)d_result[0] - d_data.begin();
  
  ASSERT_EQUAL(h_size, d_size);

  
  ASSERT_EQUAL(h_result, d_result);
}


template<typename T>
void TestRemoveCopyIfDeviceSeq(const size_t n)
{
  TestRemoveCopyIfDevice<T>(thrust::seq, n);
}
DECLARE_VARIABLE_UNITTEST(TestRemoveCopyIfDeviceSeq);


template<typename T>
void TestRemoveCopyIfDeviceDevice(const size_t n)
{
  TestRemoveCopyIfDevice<T>(thrust::device, n);
}
DECLARE_VARIABLE_UNITTEST(TestRemoveCopyIfDeviceDevice);


template<typename T, typename ExecutionPolicy>
void TestRemoveCopyIfStencilDevice(ExecutionPolicy exec, const size_t n)
{
  thrust::host_vector<T>   h_data = unittest::random_samples<T>(n);
  thrust::device_vector<T> d_data = h_data;
