From 50f28c9e252c342dd519163456e9b22301639099 Mon Sep 17 00:00:00 2001
From: Vladislav Vinogradov <vlad.vinogradov@itseez.com>
Date: Wed, 13 Feb 2013 15:57:55 +0400
Subject: [PATCH] added Warp::reduce function

---
 .../gpu/include/opencv2/gpu/device/warp.hpp   | 21 ++++++++++++++++++-
 1 file changed, 20 insertions(+), 1 deletion(-)

diff --git a/modules/gpu/include/opencv2/gpu/device/warp.hpp b/modules/gpu/include/opencv2/gpu/device/warp.hpp
index d4b0b8d8f..0f1dc794a 100644
--- a/modules/gpu/include/opencv2/gpu/device/warp.hpp
+++ b/modules/gpu/include/opencv2/gpu/device/warp.hpp
@@ -97,6 +97,25 @@ namespace cv { namespace gpu { namespace device
             return out;
         }
 
+        template <class T, class BinOp>
+        static __device__ __forceinline__ T reduce(volatile T *ptr, BinOp op)
+        {
+            const unsigned int lane = laneId();
+
+            if (lane < 16)
+            {
+                T partial = ptr[lane];
+
+                ptr[lane] = partial = op(partial, ptr[lane + 16]);
+                ptr[lane] = partial = op(partial, ptr[lane + 8]);
+                ptr[lane] = partial = op(partial, ptr[lane + 4]);
+                ptr[lane] = partial = op(partial, ptr[lane + 2]);
+                ptr[lane] = partial = op(partial, ptr[lane + 1]);
+            }
+
+            return *ptr;
+        }
+
         template<typename OutIt, typename T>
         static __device__ __forceinline__ void yota(OutIt beg, OutIt end, T value)
         {
@@ -109,4 +128,4 @@ namespace cv { namespace gpu { namespace device
     };
 }}} // namespace cv { namespace gpu { namespace device
 
-#endif /* __OPENCV_GPU_DEVICE_WARP_HPP__ */
\ No newline at end of file
+#endif /* __OPENCV_GPU_DEVICE_WARP_HPP__ */