- thrust::copy() called from dvec::copy() for gpairs invoked a GPU kernel instead of
cudaMemcpy()
- this resulted in illegal memory access if the GPU running the kernel could not access
the data being copied
- new version of dvec::copy() for thrust::device_ptr iterators calls cudaMemcpy(),
avoiding the problem.
