   
   thrust::for_each_n(h_input.begin(), h_input.size(), h_f);
   
-  test_for_each_n_device_seq_kernel<<<1,1>>>(d_input.begin(), d_input.size(), d_f);
+  for_each_n_kernel<<<1,1>>>(thrust::seq, d_input.begin(), d_input.size(), d_f);
   
   ASSERT_EQUAL(h_output, d_output);
 }
 DECLARE_VARIABLE_UNITTEST(TestForEachNDeviceSeq);
 
+
+template<typename T>
+void TestForEachNDeviceDevice(const size_t n)
+{
+  const size_t output_size = std::min((size_t) 10, 2 * n);
+  
+  thrust::host_vector<T> h_input = unittest::random_integers<T>(n);
+  
+  for(size_t i = 0; i < n; i++)
+    h_input[i] =  ((size_t) h_input[i]) % output_size;
+  
+  thrust::device_vector<T> d_input = h_input;
+  
+  thrust::host_vector<T>   h_output(output_size, (T) 0);
+  thrust::device_vector<T> d_output(output_size, (T) 0);
+  
+  mark_present_for_each<T> h_f;
+  mark_present_for_each<T> d_f;
+  h_f.ptr = &h_output[0];
+  d_f.ptr = (&d_output[0]).get();
+  
+  thrust::for_each_n(h_input.begin(), h_input.size(), h_f);
+  
+  for_each_n_kernel<<<1,1>>>(thrust::device, d_input.begin(), d_input.size(), d_f);
+  
+  ASSERT_EQUAL(h_output, d_output);
+}
+DECLARE_VARIABLE_UNITTEST(TestForEachNDeviceDevice);
+

};


template<typename Iterator, typename Function>
__global__ void test_for_each_device_seq_kernel(Iterator first, Iterator last, Function f)
{
  thrust::for_each(thrust::seq, first, last, f);
}



  
  thrust::for_each(h_input.begin(), h_input.end(), h_f);
  
  for_each_kernel<<<1,1>>>(thrust::seq, d_input.begin(), d_input.end(), d_f);
  
  ASSERT_EQUAL(h_output, d_output);
}
DECLARE_VARIABLE_UNITTEST(TestForEachDeviceSeq);


template<typename T>
void TestForEachDeviceDevice(const size_t n)
{
  const size_t output_size = std::min((size_t) 10, 2 * n);
  
  thrust::host_vector<T> h_input = unittest::random_integers<T>(n);
  
  for(size_t i = 0; i < n; i++)
    h_input[i] =  ((size_t) h_input[i]) % output_size;
  
  thrust::device_vector<T> d_input = h_input;
  
  thrust::host_vector<T>   h_output(output_size, (T) 0);
  thrust::device_vector<T> d_output(output_size, (T) 0);
  
  mark_present_for_each<T> h_f;
  mark_present_for_each<T> d_f;
  h_f.ptr = &h_output[0];
  d_f.ptr = (&d_output[0]).get();
  
  thrust::for_each(h_input.begin(), h_input.end(), h_f);
  
  for_each_kernel<<<1,1>>>(thrust::device, d_input.begin(), d_input.end(), d_f);
  
  ASSERT_EQUAL(h_output, d_output);
}
DECLARE_VARIABLE_UNITTEST(TestForEachDeviceDevice);


template<typename ExecutionPolicy, typename Iterator, typename Size, typename Function>
__global__
void for_each_n_kernel(ExecutionPolicy exec, Iterator first, Size n, Function f)
{
  thrust::for_each_n(exec, first, n, f);
}


