 #include <thrust/execution_policy.h>
 
 
-template<typename Iterator, typename Function>
+template<typename ExecutionPolicy, typename Iterator, typename Function>
 __global__
-void tabulate_kernel(Iterator first, Iterator last, Function f)
+void tabulate_kernel(ExecutionPolicy exec, Iterator first, Iterator last, Function f)
 {
-  thrust::tabulate(thrust::seq, first, last, f);
+  thrust::tabulate(exec, first, last, f);
 }
 
 
-void TestTabulateDeviceSeq()
+template<typename ExecutionPolicy>
+void TestTabulateDevice(ExecutionPolicy exec)
 {
   typedef thrust::device_vector<int> Vector;
   using namespace thrust::placeholders;

  ASSERT_EQUAL(v[3], 3);
  ASSERT_EQUAL(v[4], 4);

  tabulate_kernel<<<1,1>>>(v.begin(), v.end(), -_1);

  ASSERT_EQUAL(v[0],  0);
  ASSERT_EQUAL(v[1], -1);

  
  Vector v(5);

  tabulate_kernel<<<1,1>>>(exec, v.begin(), v.end(), thrust::identity<T>());

  ASSERT_EQUAL(v[0], 0);
  ASSERT_EQUAL(v[1], 1);
