 };
 
 
-template<typename T>
-void TestGatherIfDeviceSeq(const size_t n)
+template<typename T, typename ExecutionPolicy>
+void TestGatherIfDevice(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(h_map.begin(), h_map.end(), h_source.begin(), h_output.begin());
  gather_kernel<<<1,1>>>(d_map.begin(), d_map.end(), d_source.begin(), d_output.begin());
  
  ASSERT_EQUAL(h_output, d_output);
}
DECLARE_VARIABLE_UNITTEST(TestGatherDeviceSeq);


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



#include <thrust/execution_policy.h>


template<typename ExecutionPolicy, typename Iterator1, typename Iterator2, typename Iterator3>
__global__
void gather_kernel(ExecutionPolicy exec, Iterator1 map_first, Iterator1 map_last, Iterator2 elements_first, Iterator3 result)
{
  thrust::gather(exec, map_first, map_last, elements_first, result);
}


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);
  
