mirror/opencv
1
0
Fork 0
mirror of https://github.com/opencv/opencv.git synced 2025-01-18 06:03:15 +08:00
Code Issues Actions Packages Projects Releases Wiki Activity

added Warp::reduce function

Browse Source
This commit is contained in:
Vladislav Vinogradov 2013-02-13 15:57:55 +04:00
parent 36e42084f0
commit 50f28c9e25
1 changed files with 20 additions and 1 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

19
modules/gpu/include/opencv2/gpu/device/warp.hpp
View File

@ -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)
{