 #include <thrust/execution_policy.h>
 
 
-template<typename Iterator1, typename Iterator2, typename Function, typename Iterator3>
-__global__
-void transform_kernel(Iterator1 first, Iterator1 last, Iterator2 result1, Function f, Iterator3 result2)
-{
-  *result2 = thrust::transform(thrust::seq, first, last, result1, f);
-}
-
-
-void TestTransformUnaryDeviceSeq()
+template<typename ExecutionPolicy, typename Iterator1, typename Iterator2, typename Function, typename Iterator3>
+__global__
+void transform_kernel(ExecutionPolicy exec, Iterator1 first, Iterator1 last, Iterator2 result1, Function f, Iterator3 result2)
+{
+  *result2 = thrust::transform(exec, first, last, result1, f);
+}
+
+
+template<typename ExecutionPolicy>
+void TestTransformUnaryDevice(ExecutionPolicy exec)
 {
   typedef thrust::device_vector<int> Vector;
   typedef typename Vector::value_type T;

  ASSERT_EQUAL(iter - output.begin(), input.size());
  ASSERT_EQUAL(output, result);
}
DECLARE_UNITTEST(TestTransformIfUnaryNoStencilDeviceSeq);


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


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


  thrust::device_vector<typename Vector::iterator> iter_vec(1);
  
  transform_if_kernel<<<1,1>>>(exec,
                               input.begin(), input.end(),
                               output.begin(),
                               thrust::negate<T>(),
                               thrust::identity<T>(),
