     ASSERT_EQUAL(ref, data);
   }
 }
+
+
+void TestStablePartitionDeviceSeq()
+{
+  TestStablePartitionDevice(thrust::seq);
+}
 DECLARE_UNITTEST(TestStablePartitionDeviceSeq);
 
 
-template<typename Iterator1, typename Iterator2, typename Predicate, typename Iterator3, typename Iterator4>
-__global__
-void stable_partition_kernel(Iterator1 first, Iterator1 last, Iterator2 stencil_first, Predicate pred, Iterator3 result, Iterator4 is_supported)
+void TestStablePartitionDeviceDevice()
+{
+  TestStablePartitionDevice(thrust::device);
+}
+DECLARE_UNITTEST(TestStablePartitionDeviceDevice);
+
+
+template<typename ExecutionPolicy, typename Iterator1, typename Iterator2, typename Predicate, typename Iterator3, typename Iterator4>
+__global__
+void stable_partition_kernel(ExecutionPolicy exec, Iterator1 first, Iterator1 last, Iterator2 stencil_first, Predicate pred, Iterator3 result, Iterator4 is_supported)
 {
 #if (__CUDA_ARCH__ >= 200)
   *is_supported = true;
-  *result = thrust::stable_partition(thrust::seq, first, last, stencil_first, pred);
+  *result = thrust::stable_partition(exec, first, last, stencil_first, pred);
 #else
   *is_supported = false;
 #endif
 }
 
 
-void TestStablePartitionStencilDeviceSeq()
+template<typename ExecutionPolicy>
+void TestStablePartitionStencilDevice(ExecutionPolicy exec)
 {
   typedef int T;
   typedef typename thrust::device_vector<T>::iterator iterator;

  ASSERT_EQUAL(true_ref, true_results);
  ASSERT_EQUAL(false_ref, false_results);
}
DECLARE_UNITTEST(TestPartitionCopyDeviceSeq);


template<typename Iterator1, typename Iterator2, typename Iterator3, typename Iterator4, typename Predicate, typename Iterator5>
__global__
void partition_copy_kernel(Iterator1 first, Iterator1 last, Iterator2 stencil_first, Iterator3 true_result, Iterator4 false_result, Predicate pred, Iterator5 result)
{
  *result = thrust::partition_copy(thrust::seq, first, last, stencil_first, true_result, false_result, pred);
}


template<typename Vector>
void TestPartitionCopyStencilDeviceSeq()
{
  typedef int T;
  

  thrust::device_vector<iterator> result(1);
  thrust::device_vector<bool> is_supported(1);
  
  stable_partition_kernel<<<1,1>>>(exec, data.begin(), data.end(), is_even<T>(), result.begin(), is_supported.begin());
  
  if(is_supported[0])
  {
