 };
 
 
-template<typename Iterator, typename T1, typename T2>
-__global__
-void replace_kernel(Iterator first, Iterator last, T1 old_value, T2 new_value)
-{
-  thrust::replace(thrust::seq, first, last, old_value, new_value);
-}
-
-
-template<typename T>
-void TestReplaceDeviceSeq(const size_t n)
+template<typename ExecutionPolicy, typename Iterator, typename T1, typename T2>
+__global__
+void replace_kernel(ExecutionPolicy exec, Iterator first, Iterator last, T1 old_value, T2 new_value)
+{
+  thrust::replace(exec, first, last, old_value, new_value);
+}
+
+
+template<typename T, typename ExecutionPolicy>
+void TestReplaceDevice(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<T> d_dest(n);
  
  thrust::replace_copy_if(h_data.begin(), h_data.end(), h_dest.begin(), less_than_five<T>(), 0);
  replace_copy_if_kernel<<<1,1>>>(d_data.begin(), d_data.end(), d_dest.begin(), less_than_five<T>(), 0);
  
  ASSERT_ALMOST_EQUAL(h_data, d_data);
  ASSERT_ALMOST_EQUAL(h_dest, d_dest);
}
DECLARE_VARIABLE_UNITTEST(TestReplaceCopyIfDeviceSeq);


template<typename Iterator1, typename Iterator2, typename Iterator3, typename Predicate, typename T>
__global__
void replace_copy_if_kernel(Iterator1 first, Iterator1 last, Iterator2 stencil_first, Iterator3 result, Predicate pred, T new_value)
{
  thrust::replace_copy_if(thrust::seq, first, last, stencil_first, result, pred, new_value);
}


template<typename T>
void TestReplaceCopyIfStencilDeviceSeq(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<T> d_dest(n);
  
  thrust::replace_copy(h_data.begin(), h_data.end(), h_dest.begin(), old_value, new_value);
  replace_copy_kernel<<<1,1>>>(exec, d_data.begin(), d_data.end(), d_dest.begin(), old_value, new_value);
  
  ASSERT_ALMOST_EQUAL(h_data, d_data);
  ASSERT_ALMOST_EQUAL(h_dest, d_dest);
}

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

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


template<typename ExecutionPolicy, typename Iterator, typename Predicate, typename T>
__global__
void replace_if_kernel(ExecutionPolicy exec, Iterator first, Iterator last, Predicate pred, T new_value)
{
  thrust::replace_if(exec, first, last, pred, new_value);
}


template<typename T, typename ExecutionPolicy>
void TestReplaceIfDevice(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::replace_if(h_data.begin(), h_data.end(), less_than_five<T>(), (T) 0);
  replace_if_kernel<<<1,1>>>(exec, d_data.begin(), d_data.end(), less_than_five<T>(), (T) 0);
  
  ASSERT_ALMOST_EQUAL(h_data, d_data);
}

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

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


template<typename ExecutionPolicy, typename Iterator1, typename Iterator2, typename Predicate, typename T>
__global__
void replace_if_kernel(ExecutionPolicy exec, Iterator1 first, Iterator1 last, Iterator2 stencil_first, Predicate pred, T new_value)
{
  thrust::replace_if(exec, first, last, stencil_first, pred, new_value);
}


template<typename T, typename ExecutionPolicy>
void TestReplaceIfStencilDevice(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;
