 #include <thrust/execution_policy.h>
 
 
-template<typename Iterator, typename T, typename Iterator2>
+template<typename ExecutionPolicy, typename Iterator, typename T, typename Iterator2>
 __global__
-void count_kernel(Iterator first, Iterator last, T value, Iterator2 result)
+void count_kernel(ExecutionPolicy exec, Iterator first, Iterator last, T value, Iterator2 result)
 {
-  *result = thrust::count(thrust::seq, first, last, value);
+  *result = thrust::count(exec, first, last, value);
 }
 
 
-template<typename T>
-void TestCountDeviceSeq(const size_t n)
+template<typename T, typename ExecutionPolicy>
+void TestCountDevice(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<size_t> d_result(1);
  
  size_t h_result = thrust::count_if(h_data.begin(), h_data.end(), greater_than_five<T>());
  count_if_kernel<<<1,1>>>(d_data.begin(), d_data.end(), greater_than_five<T>(), d_result.begin());
  
  ASSERT_EQUAL(h_result, d_result[0]);
}
DECLARE_VARIABLE_UNITTEST(TestCountIfDeviceSeq);


};


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