From 24c2e41287d7fb73f99fe5e71eca1eb456a553ba Mon Sep 17 00:00:00 2001
From: "Andrew V. Adinetz" <adinetz@gmail.com>
Date: Thu, 1 Feb 2018 04:54:46 +0100
Subject: [PATCH] Fixed the bug with illegal memory access in
 test_large_sizes.py with 4 GPUs. (#3068)

- 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.
---
 src/common/device_helpers.cuh | 10 ++++++++++
 1 file changed, 10 insertions(+)

diff --git a/src/common/device_helpers.cuh b/src/common/device_helpers.cuh
index 2d7c602d3..c10c6a4eb 100644
--- a/src/common/device_helpers.cuh
+++ b/src/common/device_helpers.cuh
@@ -369,6 +369,16 @@ class dvec {
     }
     thrust::copy(begin, end, this->tbegin());
   }
+
+  void copy(thrust::device_ptr<T> begin, thrust::device_ptr<T> end) {
+    safe_cuda(cudaSetDevice(this->device_idx()));
+    if (end - begin != size()) {
+      throw std::runtime_error(
+                               "Cannot copy assign vector to dvec, sizes are different");
+    }
+    safe_cuda(cudaMemcpy(this->data(), begin.get(),
+                         size() * sizeof(T), cudaMemcpyDefault));
+  }
 };
 
 /**