From c80d51ccb3942be0884cbecd74fe9436f146171e Mon Sep 17 00:00:00 2001
From: Thejaswi <rao.thejaswi@gmail.com>
Date: Fri, 4 May 2018 06:44:08 +0530
Subject: [PATCH] Fix issue #3264, accuracy issues on k80 GPUs. (#3293)

---
 src/tree/updater_gpu_hist.cu | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/src/tree/updater_gpu_hist.cu b/src/tree/updater_gpu_hist.cu
index 93309f3db..7fc74bcd7 100644
--- a/src/tree/updater_gpu_hist.cu
+++ b/src/tree/updater_gpu_hist.cu
@@ -37,9 +37,9 @@ __device__ GradientPairSumT ReduceFeature(const GradientPairSumT* begin,
     bool thread_active = itr + threadIdx.x < end;
     // Scan histogram
     GradientPairSumT bin = thread_active ? *(itr + threadIdx.x) : GradientPairSumT();
-
-    local_sum += ReduceT(temp_storage->sum_reduce).Reduce(bin, cub::Sum());
+    local_sum += bin;
   }
+  local_sum = ReduceT(temp_storage->sum_reduce).Reduce(local_sum, cub::Sum());
 
   if (threadIdx.x == 0) {
     shared_sum = local_sum;