   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>());
+  gather_if_kernel<<<1,1>>>(exec, 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);
 }
+
+template<typename T>
+void TestGatherIfDeviceSeq(const size_t n)
+{
+  TestGatherIfDevice<T>(thrust::seq, n);
+}
 DECLARE_VARIABLE_UNITTEST(TestGatherIfDeviceSeq);
 
+template<typename T>
+void TestGatherIfDeviceDevice(const size_t n)
+{
+  TestGatherIfDevice<T>(thrust::device, n);
+}
+DECLARE_VARIABLE_UNITTEST(TestGatherIfDeviceDevice);
+

#include <thrust/execution_policy.h>


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


template<typename T>
void TestGatherDeviceSeq(const size_t n)
{
  const size_t source_size = std::min((size_t) 10, 2 * 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);
  
