 #include <thrust/execution_policy.h>
 
 
-template<typename Iterator1, typename Iterator2, typename Iterator3>
+template<typename ExecutionPolicy, typename Iterator1, typename Iterator2, typename Iterator3>
 __global__
-void gather_kernel(Iterator1 map_first, Iterator1 map_last, Iterator2 elements_first, Iterator3 result)
+void gather_kernel(ExecutionPolicy exec, Iterator1 map_first, Iterator1 map_last, Iterator2 elements_first, Iterator3 result)
 {
-  thrust::gather(thrust::seq, map_first, map_last, elements_first, result);
+  thrust::gather(exec, map_first, map_last, elements_first, result);
 }
 
 
-template<typename T>
-void TestGatherDeviceSeq(const size_t n)
+template<typename T, typename ExecutionPolicy>
+void TestGatherDevice(ExecutionPolicy exec, const size_t n)
 {
   const size_t source_size = std::min((size_t) 10, 2 * n);
   

  thrust::device_vector<T> d_output(n);
  
  thrust::gather_if(h_map.begin(), h_map.end(), h_stencil.begin(), h_source.begin(), h_output.begin(), is_even_gather_if<unsigned int>());
  gather_if_kernel<<<1,1>>>(d_map.begin(), d_map.end(), d_stencil.begin(), d_source.begin(), d_output.begin(), is_even_gather_if<unsigned int>());
  
  ASSERT_EQUAL(h_output, d_output);
}
DECLARE_VARIABLE_UNITTEST(TestGatherIfDeviceSeq);


  thrust::device_vector<T> d_output(n);
  
  thrust::gather(h_map.begin(), h_map.end(), h_source.begin(), h_output.begin());
  gather_kernel<<<1,1>>>(exec, d_map.begin(), d_map.end(), d_source.begin(), d_output.begin());
  
  ASSERT_EQUAL(h_output, d_output);
}

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

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


template<typename ExecutionPolicy, typename Iterator1, typename Iterator2, typename Iterator3, typename Iterator4, typename Predicate>
__global__
void gather_if_kernel(ExecutionPolicy exec, Iterator1 map_first, Iterator1 map_last, Iterator2 stencil_first, Iterator3 elements_first, Iterator4 result, Predicate pred)
{
  thrust::gather_if(exec, map_first, map_last, stencil_first, elements_first, result, pred);
}


