   thrust::device_vector<T> d_output(output_size, (T) 0);
   
   thrust::scatter_if(h_input.begin(), h_input.end(), h_map.begin(), h_map.begin(), h_output.begin(), is_even_scatter_if<unsigned int>());
-  scatter_if_kernel<<<1,1>>>(d_input.begin(), d_input.end(), d_map.begin(), d_map.begin(), d_output.begin(), is_even_scatter_if<unsigned int>());
+  scatter_if_kernel<<<1,1>>>(exec, d_input.begin(), d_input.end(), d_map.begin(), d_map.begin(), d_output.begin(), is_even_scatter_if<unsigned int>());
   
   ASSERT_EQUAL(h_output, d_output);
 }
+
+template<typename T>
+void TestScatterIfDeviceSeq(const size_t n)
+{
+  TestScatterIfDevice<T>(thrust::seq, n);
+}
 DECLARE_VARIABLE_UNITTEST(TestScatterIfDeviceSeq);
 
+template<typename T>
+void TestScatterIfDeviceDevice(const size_t n)
+{
+  TestScatterIfDevice<T>(thrust::device, n);
+}
+DECLARE_VARIABLE_UNITTEST(TestScatterIfDeviceDevice);
+

};


template<typename T>
void TestScatterIfDeviceSeq(const size_t n)
{
  const size_t output_size = std::min((size_t) 10, 2 * n);
  

#include <thrust/execution_policy.h>


template<typename ExecutionPolicy, typename Iterator1, typename Iterator2, typename Iterator3>
__global__
void scatter_kernel(ExecutionPolicy exec, Iterator1 first, Iterator1 last, Iterator2 map_first, Iterator3 result)
{
  thrust::scatter(exec, first, last, map_first, result);
}


template<typename T, typename ExecutionPolicy>
void TestScatterDevice(ExecutionPolicy exec, const size_t n)
{
  const size_t output_size = std::min((size_t) 10, 2 * n);
  
