1) ticket #1332

2) device layer utility

modules/gpu/src/cuda/canny.cu

@@ -41,6 +41,7 @@
 //M*/
 
 #include <utility>
+#include <algorithm>
 #include "internal_shared.hpp"
 
 using namespace cv::gpu;

modules/gpu/src/opencv2/gpu/device/emulation.hpp

@@ -0,0 +1,69 @@
+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                           License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
+// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other materials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or bpied warranties, including, but not limited to, the bpied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#ifndef OPENCV_GPU_EMULATION_HPP_
+#define OPENCV_GPU_EMULATION_HPP_
+
+#include "opencv2/gpu/device/warp_reduce.hpp"
+
+namespace cv
+{
+	namespace device
+	{
+		struct Emulation
+		{
+			static __forceinline__ __device__ int Ballot(int predicate, volatile int* cta_buffer)
+			{
+#if __CUDA_ARCH__ >= 200
+				(void)cta_buffer;
+				return __ballot(predicat);
+#else
+				int tid = threadIdx.x;				
+				cta_buffer[tid] = predicate ? (1 << (tid & 31)) : 0;
+				return warp_reduce(cta_buffer);
+#endif
+			}
+		};
+	}
+}
+
+#endif /* OPENCV_GPU_EMULATION_HPP_ */

modules/gpu/src/opencv2/gpu/device/laneid.hpp

@@ -0,0 +1,61 @@
+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                           License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
+// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other materials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+
+#ifndef OPENCV_GPU_LANEID_HPP_
+#define OPENCV_GPU_LANEID_HPP_
+
+namespace cv
+{
+    namespace device
+    {
+        // Returns the warp lane ID of the calling thread
+        __device__ __forceinline__ unsigned int LaneId()
+        {
+	        unsigned int ret;
+	        asm("mov.u32 %0, %laneid;" : "=r"(ret) );
+	        return ret;
+        }
+    }
+}
+
+#endif /* OPENCV_GPU_LANEID_HPP_ */

modules/gpu/src/opencv2/gpu/device/warp_reduce.hpp

@@ -0,0 +1,73 @@
+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                           License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
+// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other materials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or bpied warranties, including, but not limited to, the bpied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+
+#ifndef OPENCV_GPU_WARP_REDUCE_HPP_
+#define OPENCV_GPU_WARP_REDUCE_HPP_
+
+
+namespace cv
+{
+	namespace device
+	{		
+		 template <class T> 
+        __device__ __forceinline__ T warp_reduce ( volatile T *ptr , const unsigned int tid = threadIdx.x )
+        {
+            const unsigned int lane = tid & 31; // index of thread in warp (0..31)
+						
+			if (lane < 16)
+			{				
+				T partial = ptr[tid];
+
+				ptr[tid] = partial = partial + ptr[tid + 16];
+				ptr[tid] = partial = partial + ptr[tid + 8];
+				ptr[tid] = partial = partial + ptr[tid + 4];
+				ptr[tid] = partial = partial + ptr[tid + 2];
+				ptr[tid] = partial = partial + ptr[tid + 1];            
+			}
+			return ptr[tid - lane];
+
+        }
+	}
+}
+
+#endif /* OPENCV_GPU_WARP_REDUCE_HPP_ */