   
   // inplace inclusive scan
   input = input_copy;
-  transform_inclusive_scan_kernel<<<1,1>>>(input.begin(), input.end(), input.begin(), thrust::negate<T>(), thrust::plus<T>(), iter_vec.begin());
+  transform_inclusive_scan_kernel<<<1,1>>>(exec, input.begin(), input.end(), input.begin(), thrust::negate<T>(), thrust::plus<T>(), iter_vec.begin());
   iter = iter_vec[0];
   ref[0] = -1; ref[1] = -4; ref[2] = -2; ref[3] = -6; ref[4] = -1;
   ASSERT_EQUAL(iter - input.begin(), input.size());

  thrust::device_vector<typename Vector::iterator> iter_vec(1);
  
  // inclusive scan
  transform_inclusive_scan_kernel<<<1,1>>>(input.begin(), input.end(), output.begin(), thrust::negate<T>(), thrust::plus<T>(), iter_vec.begin());
  iter = iter_vec[0];
  ref[0] = -1; ref[1] = -4; ref[2] = -2; ref[3] = -6; ref[4] = -1;
  ASSERT_EQUAL(iter - output.begin(), input.size());

  ASSERT_EQUAL(ref, output);
  
  // exclusive scan with 0 init
  transform_exclusive_scan_kernel<<<1,1>>>(exec, input.begin(), input.end(), output.begin(), thrust::negate<T>(), 0, thrust::plus<T>(), iter_vec.begin());
  ref[0] = 0; ref[1] = -1; ref[2] = -4; ref[3] = -2; ref[4] = -6;
  ASSERT_EQUAL(iter - output.begin(), input.size());
  ASSERT_EQUAL(input,  input_copy);
  ASSERT_EQUAL(ref, output);
  
  // exclusive scan with nonzero init
  transform_exclusive_scan_kernel<<<1,1>>>(exec, input.begin(), input.end(), output.begin(), thrust::negate<T>(), 3, thrust::plus<T>(), iter_vec.begin());
  iter = iter_vec[0];
  ref[0] = 3; ref[1] = 2; ref[2] = -1; ref[3] = 1; ref[4] = -3;
  ASSERT_EQUAL(iter - output.begin(), input.size());
