Infer the directionality of trivial_copy_n dynamically

Because the user may have explicitly passed an execution policy such as thrust::cuda::par.on(stream), trivial_copy_n could be cross-space. Dynamically infer the directionality of the copy using cudaMemcpyDefault.

Fixes #651

thrust/system/cuda/detail/trivial_copy.inl

@@ -86,7 +86,12 @@ void trivial_copy_n(execution_policy<DerivedPolicy> &exec,
   void *dst = thrust::raw_pointer_cast(&*result);
   const void *src = thrust::raw_pointer_cast(&*first);
 
-  trivial_copy_detail::checked_cudaMemcpyAsync(dst, src, n * sizeof(T), cudaMemcpyDeviceToDevice, stream(thrust::detail::derived_cast(exec)));
+  // since the user may have given thrust::cuda::par to thrust::copy explicitly,
+  // this copy may be a cross-space copy that has bypassed system dispatch
+  // we need to have cudaMemcpyAsync figure out the directionality of the copy dynamically
+  // using cudaMemcpyDefault
+
+  trivial_copy_detail::checked_cudaMemcpyAsync(dst, src, n * sizeof(T), cudaMemcpyDefault, stream(thrust::detail::derived_cast(exec)));
 #else
   thrust::transform(exec, first, first + n, result, thrust::identity<T>());
 #endif