generic-library/opencv
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

experimental kernels for cuda

Browse Source
This commit is contained in:
Anatoly Baksheev 2011-08-10 09:51:36 +00:00
parent 50d1d711de
commit edcfa64d99
1 changed files with 16 additions and 3 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

19
modules/gpu/src/opencv2/gpu/device/kernels.hpp
View File

@ -63,7 +63,7 @@ namespace cv
struct Warp 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); } static __forceinline__ __device__ int SHIFT() { return threadIdx.x & (warpSize - 1); }
}; };
@ -77,8 +77,8 @@ namespace cv
out[idx] = in[idx]; out[idx] = in[idx];
} }
template <class Worker, typename ForwardIterator, typename ForwardIterator> template <class Worker, typename InIter, typename OutIter>
__forceinline__ __device__ void Copy(ForwardIterator beg, ForwardIterator end, OutIter out) __forceinline__ __device__ void Copy(InIter beg, InIter end, OutIter out)
{ {
int STRIDE = Worker::STRIDE(); int STRIDE = Worker::STRIDE();
int SHIFT = Worker::SHIFT(); int SHIFT = Worker::SHIFT();
@ -103,6 +103,19 @@ namespace cv
for (; idx < length; idx += STRIDE, cur += STRIDE) for (; idx < length; idx += STRIDE, cur += STRIDE)
out[idx] = cur; 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;
}
} }
} }
} }