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


  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());
  iter = iter_vec[0];
  
  ASSERT_EQUAL(iter - output.begin(), input.size());
  ASSERT_EQUAL(output, result);
}
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()
{
  typedef thrust::device_vector<int> Vector;
  typedef typename Vector::value_type T;

  output[0]  =  1; output[1]  =  2; output[2]  =  3; 
  stencil[0] =  1; stencil[1] =  0; stencil[2] =  1;
  result[0]  = -1; result[1]  =  2; result[2]  = -3;

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

  iter = iter_vec[0];
  
  ASSERT_EQUAL(iter - output.begin(), input.size());
  ASSERT_EQUAL(output, result);
}

void TestTransformIfUnaryDeviceSeq()
{
  TestTransformIfUnaryDevice(thrust::seq);
}
DECLARE_UNITTEST(TestTransformIfUnaryDeviceSeq);

void TestTransformIfUnaryDeviceDevice()
{
  TestTransformIfUnaryDevice(thrust::device);
}
DECLARE_UNITTEST(TestTransformIfUnaryDeviceDevice);


template<typename ExecutionPolicy, typename Iterator1, typename Iterator2, typename Iterator3, typename Function, typename Iterator4>
__global__
void transform_kernel(ExecutionPolicy exec, Iterator1 first1, Iterator1 last1, Iterator2 first2, Iterator3 result1, Function f, Iterator4 result2)
{
  *result2 = thrust::transform(exec, first1, last1, first2, result1, f);
}


template<typename ExecutionPolicy>
void TestTransformBinaryDevice(ExecutionPolicy exec)
{
  typedef thrust::device_vector<int> Vector;
  typedef typename Vector::value_type T;
