From edcfa64d9948d73e12b5ca2191723d0b3f7eb811 Mon Sep 17 00:00:00 2001
From: Anatoly Baksheev <no@email>
Date: Wed, 10 Aug 2011 09:51:36 +0000
Subject: [PATCH] experimental kernels for cuda

---
 .../gpu/src/opencv2/gpu/device/kernels.hpp    | 19 ++++++++++++++++---
 1 file changed, 16 insertions(+), 3 deletions(-)

diff --git a/modules/gpu/src/opencv2/gpu/device/kernels.hpp b/modules/gpu/src/opencv2/gpu/device/kernels.hpp
index 5298befbc..5f3f04161 100644
--- a/modules/gpu/src/opencv2/gpu/device/kernels.hpp
+++ b/modules/gpu/src/opencv2/gpu/device/kernels.hpp
@@ -63,7 +63,7 @@ namespace cv
 
 			struct Warp
 			{
-				static __forceinline__ __device__ int STRIDE() { return warpSize;            
+				static __forceinline__ __device__ int STRIDE() { return warpSize };            
 				static __forceinline__ __device__ int SHIFT()  { return threadIdx.x & (warpSize - 1); }			
 			};
 
@@ -77,8 +77,8 @@ namespace cv
 	                out[idx] = in[idx];
 			}
 
-			template <class Worker, typename ForwardIterator, typename ForwardIterator>
-			__forceinline__ __device__ void Copy(ForwardIterator beg, ForwardIterator end, OutIter out)
+			template <class Worker, typename InIter, typename OutIter>
+			__forceinline__ __device__ void Copy(InIter beg, InIter end, OutIter out)
 			{
 				int STRIDE = Worker::STRIDE();
 				int SHIFT  = Worker::SHIFT();
@@ -103,6 +103,19 @@ namespace cv
 				for (; idx < length; idx += STRIDE, cur += STRIDE)
 					out[idx] = cur;				
 			}
+
+			template <class Worker, typename OutIter>
+			__forceinline__ __device__ void Yota(OutIter beg, OutIter end, int val)
+			{
+				int STRIDE = Worker::STRIDE();
+				int SHIFT  = Worker::SHIFT();
+
+				beg += SHIFT;
+				val += SHIFT;
+				
+				for (; beg < end; beg += STRIDE, val += STRIDE)
+					*beg = val;
+			}
 		}
 	}
 }