 
   thrust::device_vector<typename Vector::iterator> iter_vec(1);
   
-  transform_kernel<<<1,1>>>(input.begin(), input.end(), output.begin(), thrust::negate<T>(), iter_vec.begin());
+  transform_kernel<<<1,1>>>(exec, input.begin(), input.end(), output.begin(), thrust::negate<T>(), iter_vec.begin());
   iter = iter_vec[0];
   
   ASSERT_EQUAL(iter - output.begin(), input.size());
   ASSERT_EQUAL(output, result);
 }
+
+void TestTransformUnaryDeviceSeq()
+{
+  TestTransformUnaryDevice(thrust::seq);
+}
 DECLARE_UNITTEST(TestTransformUnaryDeviceSeq);
 
-
-template<typename Iterator1, typename Iterator2, typename Function, typename Predicate, typename Iterator3>
-__global__
-void transform_if_kernel(Iterator1 first, Iterator1 last, Iterator2 result1, Function f, Predicate pred, Iterator3 result2)
-{
-  *result2 = thrust::transform_if(thrust::seq, first, last, result1, f, pred);
-}
-
-
-void TestTransformIfUnaryNoStencilDeviceSeq()
+void TestTransformUnaryDeviceDevice()
+{
+  TestTransformUnaryDevice(thrust::device);
+}
+DECLARE_UNITTEST(TestTransformUnaryDeviceDevice);
+
+
+template<typename ExecutionPolicy, typename Iterator1, typename Iterator2, typename Function, typename Predicate, typename Iterator3>
+__global__
+void transform_if_kernel(ExecutionPolicy exec, Iterator1 first, Iterator1 last, Iterator2 result1, Function f, Predicate pred, Iterator3 result2)
+{
+  *result2 = thrust::transform_if(exec, first, last, result1, f, pred);
+}
+
+
+template<typename ExecutionPolicy>
+void TestTransformIfUnaryNoStencilDevice(ExecutionPolicy exec)
 {
   typedef thrust::device_vector<int> Vector;
   typedef typename Vector::value_type T;


  thrust::device_vector<typename Vector::iterator> iter_vec(1);
  
  transform_kernel<<<1,1>>>(input1.begin(), input1.end(), input2.begin(), output.begin(), thrust::minus<T>(), iter_vec.begin());
  iter = iter_vec[0];
  
  ASSERT_EQUAL(iter - output.begin(), input1.size());
  ASSERT_EQUAL(output, result);
}
DECLARE_UNITTEST(TestTransformBinaryDeviceSeq);


template<typename Iterator1, typename Iterator2, typename Iterator3, typename Iterator4, typename Function, typename Predicate, typename Iterator5>
__global__
void transform_if_kernel(Iterator1 first1, Iterator1 last1, Iterator2 first2, Iterator3 stencil_first, Iterator4 result1, Function f, Predicate pred, Iterator5 result2)
{
  *result2 = thrust::transform_if(thrust::seq, first1, last1, first2, stencil_first, result1, f, pred);
}


void TestTransformIfBinaryDeviceSeq()
{
  typedef thrust::device_vector<int> Vector;
  typedef typename Vector::value_type T;

  result[0]  =  5; result[1]  =  2; result[2]  = -3;
  
  thrust::identity<T> identity;

  thrust::device_vector<typename Vector::iterator> iter_vec(1);
  
  transform_if_kernel<<<1,1>>>(exec,
                               input1.begin(), input1.end(),
                               input2.begin(),
                               stencil.begin(),
                               output.begin(),
                               thrust::minus<T>(),
                               thrust::not1(identity),
                               iter_vec.begin());
  iter = iter_vec[0];
  
  ASSERT_EQUAL(iter - output.begin(), input1.size());
  ASSERT_EQUAL(output, result);
}

void TestTransformIfBinaryDeviceSeq()
{
  TestTransformIfBinaryDevice(thrust::seq);
}
DECLARE_UNITTEST(TestTransformIfBinaryDeviceSeq);

void TestTransformIfBinaryDeviceDevice()
{
  TestTransformIfBinaryDevice(thrust::device);
}
DECLARE_UNITTEST(TestTransformIfBinaryDeviceDevice);

