From 8d66a27131286c3f86885429b24bc6d9cdf2f9d2 Mon Sep 17 00:00:00 2001
From: Tom Deakin <tom.deakin@bristol.ac.uk>
Date: Mon, 19 Dec 2016 05:08:19 -0700
Subject: [PATCH] [CUDA] If using managed memory, use device pointer for host
 reduction

---
 CUDAStream.cu | 11 +++++++++++
 1 file changed, 11 insertions(+)

diff --git a/CUDAStream.cu b/CUDAStream.cu
index 260c07e..603b0f0 100644
--- a/CUDAStream.cu
+++ b/CUDAStream.cu
@@ -245,12 +245,23 @@ T CUDAStream<T>::dot()
   dot_kernel<<<DOT_NUM_BLOCKS, TBSIZE, sizeof(T)*TBSIZE>>>(d_a, d_b, d_sum, array_size);
   check_error();
 
+#if defined(MANAGED) || defined(PAGEFAULT)
+  cudaDeviceSynchronize();
+  check_error();
+#else
   cudaMemcpy(sums, d_sum, DOT_NUM_BLOCKS*sizeof(T), cudaMemcpyDeviceToHost);
   check_error();
+#endif
 
   T sum = 0.0;
   for (int i = 0; i < DOT_NUM_BLOCKS; i++)
+  {
+#if defined(MANAGED) || defined(PAGEFAULT)
+    sum += d_sum[i];
+#else
     sum += sums[i];
+#endif
+  }
 
   return sum;
 }