 #include <thrust/execution_policy.h>
 
 
-template<typename Iterator1, typename Iterator2>
+template<typename ExecutionPolicy, typename Iterator1, typename Iterator2>
 __global__
-void uninitialized_copy_kernel(Iterator1 first, Iterator1 last, Iterator2 result)
+void uninitialized_copy_kernel(ExecutionPolicy exec, Iterator1 first, Iterator1 last, Iterator2 result)
 {
-  thrust::uninitialized_copy(thrust::seq, first, last, result);
+  thrust::uninitialized_copy(exec, first, last, result);
 }
 
 
-void TestUninitializedCopyDeviceSeq()
+template<typename ExecutionPolicy>
+void TestUninitializedCopyDevice(ExecutionPolicy exec)
 {
   typedef thrust::device_vector<int> Vector;
   typedef typename Vector::value_type T;

  
  // copy to Vector
  Vector v2(5);
  uninitialized_copy_kernel<<<1,1>>>(v1.begin(), v1.end(), v2.begin());
  ASSERT_EQUAL(v2[0], 0);
  ASSERT_EQUAL(v2[1], 1);
  ASSERT_EQUAL(v2[2], 2);
  ASSERT_EQUAL(v2[3], 3);
  ASSERT_EQUAL(v2[4], 4);
}
DECLARE_UNITTEST(TestUninitializedCopyDeviceSeq);


template<typename Iterator1, typename Size, typename Iterator2>
__global__
void uninitialized_copy_n_kernel(Iterator1 first, Size n, Iterator2 result)
{
  thrust::uninitialized_copy_n(thrust::seq, first, n, result);
}


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

  
  // copy to Vector
  Vector v2(5);
  uninitialized_copy_n_kernel<<<1,1>>>(exec, v1.begin(), v1.size(), v2.begin());
  ASSERT_EQUAL(v2[0], 0);
  ASSERT_EQUAL(v2[1], 1);
  ASSERT_EQUAL(v2[2], 2);
  ASSERT_EQUAL(v2[3], 3);
  ASSERT_EQUAL(v2[4], 4);
}


void TestUninitializedCopyNDeviceSeq()
{
  TestUninitializedCopyNDevice(thrust::seq);
}
DECLARE_UNITTEST(TestUninitializedCopyNDeviceSeq);


void TestUninitializedCopyNDeviceDevice()
{
  TestUninitializedCopyNDevice(thrust::device);
}
DECLARE_UNITTEST(TestUninitializedCopyNDeviceDevice);

