- 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.