diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp index eeec91794a..404825434b 100644 --- a/modules/gpu/include/opencv2/gpu/gpu.hpp +++ b/modules/gpu/include/opencv2/gpu/gpu.hpp @@ -60,6 +60,10 @@ #pragma GCC diagnostic ignored "-Wdeprecated-declarations" #endif +#if !defined(HAVE_TBB) + #define throw_notbb() CV_Error(CV_StsNotImplemented, "The library is compiled without TBB support") +#endif + namespace cv { namespace gpu { //////////////////////////////// CudaMem //////////////////////////////// @@ -1824,6 +1828,14 @@ public: void sparse(const GpuMat& prevImg, const GpuMat& nextImg, const GpuMat& prevPts, GpuMat& nextPts, GpuMat& status, GpuMat* err = 0); +#if !defined(HAVE_TBB) + void sparse_multi(const GpuMat&, const GpuMat&, const GpuMat&, GpuMat&, + GpuMat&, Stream&, GpuMat*) {throw_notbb();} +#else + void sparse_multi(const GpuMat& prevImg, const GpuMat& nextImg, const GpuMat& prevPts, GpuMat& nextPts, + GpuMat& status, Stream& stream, GpuMat* err = 0); +#endif + void dense(const GpuMat& prevImg, const GpuMat& nextImg, GpuMat& u, GpuMat& v, GpuMat* err = 0); void releaseMemory(); diff --git a/modules/gpu/perf/perf_video.cpp b/modules/gpu/perf/perf_video.cpp index 663b825dcc..b85fca8fd4 100644 --- a/modules/gpu/perf/perf_video.cpp +++ b/modules/gpu/perf/perf_video.cpp @@ -303,6 +303,88 @@ PERF_TEST_P(ImagePair_Gray_NPts_WinSz_Levels_Iters, Video_PyrLKOpticalFlowSparse } } +////////////////////////////////////////////////////// +// PyrLKOpticalFlowSparseMulti + +#ifdef HAVE_TBB + +DEF_PARAM_TEST(ImagePair_Gray_NPts_WinSz_Levels_Iters, pair_string, bool, int, int, int, int); + +PERF_TEST_P(ImagePair_Gray_NPts_WinSz_Levels_Iters, Video_PyrLKOpticalFlowSparseMulti, + Combine(Values(make_pair("gpu/opticalflow/frame0.png", "gpu/opticalflow/frame1.png")), + Bool(), + Values(8000), + Values(21), + Values(1, 3), + Values(1, 30))) +{ + declare.time(20.0); + + const pair_string imagePair = GET_PARAM(0); + const bool useGray = GET_PARAM(1); + const int points = GET_PARAM(2); + const int winSize = GET_PARAM(3); + const int levels = GET_PARAM(4); + const int iters = GET_PARAM(5); + + const cv::Mat frame0 = readImage(imagePair.first, useGray ? cv::IMREAD_GRAYSCALE : cv::IMREAD_COLOR); + ASSERT_FALSE(frame0.empty()); + + const cv::Mat frame1 = readImage(imagePair.second, useGray ? cv::IMREAD_GRAYSCALE : cv::IMREAD_COLOR); + ASSERT_FALSE(frame1.empty()); + + cv::Mat gray_frame; + if (useGray) + gray_frame = frame0; + else + cv::cvtColor(frame0, gray_frame, cv::COLOR_BGR2GRAY); + + cv::Mat pts; + cv::goodFeaturesToTrack(gray_frame, pts, points, 0.01, 0.0); + + if (PERF_RUN_GPU()) + { + const cv::gpu::GpuMat d_pts(pts.reshape(2, 1)); + + cv::gpu::PyrLKOpticalFlow d_pyrLK; + d_pyrLK.winSize = cv::Size(winSize, winSize); + d_pyrLK.maxLevel = levels - 1; + d_pyrLK.iters = iters; + + const cv::gpu::GpuMat d_frame0(frame0); + const cv::gpu::GpuMat d_frame1(frame1); + cv::gpu::GpuMat nextPts; + cv::gpu::GpuMat status; + + cv::gpu::Stream stream; + TEST_CYCLE() + { + d_pyrLK.sparse_multi(d_frame0, d_frame1, d_pts, nextPts, status, stream); + stream.waitForCompletion(); + } + + GPU_SANITY_CHECK(nextPts); + GPU_SANITY_CHECK(status); + } + else + { + cv::Mat nextPts; + cv::Mat status; + + TEST_CYCLE() + { + cv::calcOpticalFlowPyrLK(frame0, frame1, pts, nextPts, status, cv::noArray(), + cv::Size(winSize, winSize), levels - 1, + cv::TermCriteria(cv::TermCriteria::COUNT + cv::TermCriteria::EPS, iters, 0.01)); + } + + CPU_SANITY_CHECK(nextPts); + CPU_SANITY_CHECK(status); + } +} + +#endif // HAVE_TBB + ////////////////////////////////////////////////////// // PyrLKOpticalFlowDense diff --git a/modules/gpu/src/cuda/pyrlk.cu b/modules/gpu/src/cuda/pyrlk.cu index 8ef728b03b..518ef63fa4 100644 --- a/modules/gpu/src/cuda/pyrlk.cu +++ b/modules/gpu/src/cuda/pyrlk.cu @@ -49,6 +49,8 @@ #include "opencv2/gpu/device/vec_math.hpp" #include "opencv2/gpu/device/reduce.hpp" +#include "opencv2/core/core.hpp" + using namespace cv::gpu; using namespace cv::gpu::device; @@ -60,12 +62,54 @@ namespace pyrlk __constant__ int c_halfWin_y; __constant__ int c_iters; - texture tex_If(false, cudaFilterModeLinear, cudaAddressModeClamp); - texture tex_If4(false, cudaFilterModeLinear, cudaAddressModeClamp); - texture tex_Ib(false, cudaFilterModePoint, cudaAddressModeClamp); +#define CUDA_CONSTANTS(index) \ + __constant__ int c_winSize_x##index; \ + __constant__ int c_winSize_y##index; \ + __constant__ int c_halfWin_x##index; \ + __constant__ int c_halfWin_y##index; \ + __constant__ int c_iters##index; - texture tex_Jf(false, cudaFilterModeLinear, cudaAddressModeClamp); - texture tex_Jf4(false, cudaFilterModeLinear, cudaAddressModeClamp); + CUDA_CONSTANTS(0) + CUDA_CONSTANTS(1) + CUDA_CONSTANTS(2) + CUDA_CONSTANTS(3) + CUDA_CONSTANTS(4) + + template struct c_multi_winSize_x; + template struct c_multi_winSize_y; + template struct c_multi_halfWin_x; + template struct c_multi_halfWin_y; + template struct c_multi_iters; + +#define CUDA_CONSTANTS_ACCESSOR(index) \ + template <> struct c_multi_winSize_x \ + { static __device__ __forceinline__ int get(void){ return c_winSize_x##index;} }; \ + template <> struct c_multi_winSize_y \ + { static __device__ __forceinline__ int get(void){ return c_winSize_y##index;} }; \ + template <> struct c_multi_halfWin_x \ + { static __device__ __forceinline__ int get(void){ return c_halfWin_x##index;} }; \ + template <> struct c_multi_halfWin_y \ + { static __device__ __forceinline__ int get(void){ return c_halfWin_y##index;} }; \ + template <> struct c_multi_iters \ + { static __device__ __forceinline__ int get(void){ return c_iters##index;} }; + + CUDA_CONSTANTS_ACCESSOR(0) + CUDA_CONSTANTS_ACCESSOR(1) + CUDA_CONSTANTS_ACCESSOR(2) + CUDA_CONSTANTS_ACCESSOR(3) + CUDA_CONSTANTS_ACCESSOR(4) + + texture + tex_If(false, cudaFilterModeLinear, cudaAddressModeClamp); + texture + tex_If4(false, cudaFilterModeLinear, cudaAddressModeClamp); + texture + tex_Ib(false, cudaFilterModePoint, cudaAddressModeClamp); + + texture + tex_Jf(false, cudaFilterModeLinear, cudaAddressModeClamp); + texture + tex_Jf4(false, cudaFilterModeLinear, cudaAddressModeClamp); template struct Tex_I; template <> struct Tex_I<1> @@ -99,6 +143,57 @@ namespace pyrlk } }; + //-------------------------------------------------------------------------- + +#define CUDA_DECL_TEX_MULTI(texname, type, filtermode) \ + texture \ + texname##_multi0(false, filtermode, cudaAddressModeClamp); \ + texture \ + texname##_multi1(false, filtermode, cudaAddressModeClamp); \ + texture \ + texname##_multi2(false, filtermode, cudaAddressModeClamp); \ + texture \ + texname##_multi3(false, filtermode, cudaAddressModeClamp); \ + texture \ + texname##_multi4(false, filtermode, cudaAddressModeClamp); \ + + CUDA_DECL_TEX_MULTI(tex_If1, float, cudaFilterModeLinear) + CUDA_DECL_TEX_MULTI(tex_If4, float4, cudaFilterModeLinear) + CUDA_DECL_TEX_MULTI(tex_Ib1, uchar, cudaFilterModePoint) + CUDA_DECL_TEX_MULTI(tex_Jf1, float, cudaFilterModeLinear) + CUDA_DECL_TEX_MULTI(tex_Jf4, float4, cudaFilterModeLinear) + + template struct Tex_I_multi; + template struct Tex_J_multi; + template struct Tex_B_multi; + +#define CUDA_DECL_TEX_MULTI_ACCESS(accessorname, texname, cn, returntype) \ + template <> struct accessorname##_multi \ + { static __device__ __forceinline__ returntype read(float x, float y) \ + { return tex2D(texname##cn##_multi0, x, y); } }; \ + template <> struct accessorname##_multi \ + { static __device__ __forceinline__ returntype read(float x, float y) \ + { return tex2D(texname##cn##_multi1, x, y); } }; \ + template <> struct accessorname##_multi \ + { static __device__ __forceinline__ returntype read(float x, float y) \ + { return tex2D(texname##cn##_multi2, x, y); } }; \ + template <> struct accessorname##_multi \ + { static __device__ __forceinline__ returntype read(float x, float y) \ + { return tex2D(texname##cn##_multi3, x, y); } }; \ + template <> struct accessorname##_multi \ + { static __device__ __forceinline__ returntype read(float x, float y) \ + { return tex2D(texname##cn##_multi4, x, y); } }; + + CUDA_DECL_TEX_MULTI_ACCESS(Tex_I, tex_If, 1, float) + CUDA_DECL_TEX_MULTI_ACCESS(Tex_I, tex_If, 4, float4) + + CUDA_DECL_TEX_MULTI_ACCESS(Tex_B, tex_Ib, 1, uchar) + + CUDA_DECL_TEX_MULTI_ACCESS(Tex_J, tex_Jf, 1, float) + CUDA_DECL_TEX_MULTI_ACCESS(Tex_J, tex_Jf, 4, float4) + + //-------------------------------------------------------------------------- + __device__ __forceinline__ void accum(float& dst, float val) { dst += val; @@ -309,6 +404,200 @@ namespace pyrlk } } +#if defined(HAVE_TBB) + template + __global__ void sparseKernel_multi(const float2* prevPts, float2* nextPts, uchar* status, float* err, const int level, const int rows, const int cols) + { + #if __CUDA_ARCH__ <= 110 + const int BLOCK_SIZE = 128; + #else + const int BLOCK_SIZE = 256; + #endif + + __shared__ float smem1[BLOCK_SIZE]; + __shared__ float smem2[BLOCK_SIZE]; + __shared__ float smem3[BLOCK_SIZE]; + + const unsigned int tid = threadIdx.y * blockDim.x + threadIdx.x; + + float2 prevPt = prevPts[blockIdx.x]; + prevPt.x *= (1.0f / (1 << level)); + prevPt.y *= (1.0f / (1 << level)); + + if (prevPt.x < 0 || prevPt.x >= cols || prevPt.y < 0 || prevPt.y >= rows) + { + if (tid == 0 && level == 0) + status[blockIdx.x] = 0; + + return; + } + + prevPt.x -= c_multi_halfWin_x::get(); + prevPt.y -= c_multi_halfWin_y::get(); + + // extract the patch from the first image, compute covariation matrix of derivatives + + float A11 = 0; + float A12 = 0; + float A22 = 0; + + typedef typename TypeVec::vec_type work_type; + + work_type I_patch [PATCH_Y][PATCH_X]; + work_type dIdx_patch[PATCH_Y][PATCH_X]; + work_type dIdy_patch[PATCH_Y][PATCH_X]; + + for (int yBase = threadIdx.y, i = 0; yBase < c_multi_winSize_y::get(); yBase += blockDim.y, ++i) + { + for (int xBase = threadIdx.x, j = 0; xBase < c_multi_winSize_x::get(); xBase += blockDim.x, ++j) + { + float x = prevPt.x + xBase + 0.5f; + float y = prevPt.y + yBase + 0.5f; + + I_patch[i][j] = Tex_I_multi::read(x, y); + + // Sharr Deriv + + work_type dIdx = 3.0f * Tex_I_multi::read(x+1, y-1) + 10.0f * Tex_I_multi::read(x+1, y) + 3.0f * Tex_I_multi::read(x+1, y+1) - + (3.0f * Tex_I_multi::read(x-1, y-1) + 10.0f * Tex_I_multi::read(x-1, y) + 3.0f * Tex_I_multi::read(x-1, y+1)); + + work_type dIdy = 3.0f * Tex_I_multi::read(x-1, y+1) + 10.0f * Tex_I_multi::read(x, y+1) + 3.0f * Tex_I_multi::read(x+1, y+1) - + (3.0f * Tex_I_multi::read(x-1, y-1) + 10.0f * Tex_I_multi::read(x, y-1) + 3.0f * Tex_I_multi::read(x+1, y-1)); + + dIdx_patch[i][j] = dIdx; + dIdy_patch[i][j] = dIdy; + + accum(A11, dIdx * dIdx); + accum(A12, dIdx * dIdy); + accum(A22, dIdy * dIdy); + } + } + + reduce(smem_tuple(smem1, smem2, smem3), thrust::tie(A11, A12, A22), tid, thrust::make_tuple(plus(), plus(), plus())); + + #if __CUDA_ARCH__ >= 300 + if (tid == 0) + { + smem1[0] = A11; + smem2[0] = A12; + smem3[0] = A22; + } + #endif + + __syncthreads(); + + A11 = smem1[0]; + A12 = smem2[0]; + A22 = smem3[0]; + + float D = A11 * A22 - A12 * A12; + + if (abs_(D) < numeric_limits::epsilon()) + { + if (tid == 0 && level == 0) + status[blockIdx.x] = 0; + + return; + } + + D = 1.f / D; + + A11 *= D; + A12 *= D; + A22 *= D; + + float2 nextPt = nextPts[blockIdx.x]; + nextPt.x *= 2.f; + nextPt.y *= 2.f; + + nextPt.x -= c_multi_halfWin_x::get(); + nextPt.y -= c_multi_halfWin_y::get(); + + for (int k = 0; k < c_multi_iters::get(); ++k) + { + if (nextPt.x < -c_multi_halfWin_x::get() || nextPt.x >= cols || nextPt.y < -c_multi_halfWin_y::get() || nextPt.y >= rows) + { + if (tid == 0 && level == 0) + status[blockIdx.x] = 0; + + return; + } + + float b1 = 0; + float b2 = 0; + + for (int y = threadIdx.y, i = 0; y < c_multi_winSize_y::get(); y += blockDim.y, ++i) + { + for (int x = threadIdx.x, j = 0; x < c_multi_winSize_x::get(); x += blockDim.x, ++j) + { + work_type I_val = I_patch[i][j]; + work_type J_val = Tex_J_multi::read(nextPt.x + x + 0.5f, nextPt.y + y + 0.5f); + + work_type diff = (J_val - I_val) * 32.0f; + + accum(b1, diff * dIdx_patch[i][j]); + accum(b2, diff * dIdy_patch[i][j]); + } + } + + reduce(smem_tuple(smem1, smem2), thrust::tie(b1, b2), tid, thrust::make_tuple(plus(), plus())); + + #if __CUDA_ARCH__ >= 300 + if (tid == 0) + { + smem1[0] = b1; + smem2[0] = b2; + } + #endif + + __syncthreads(); + + b1 = smem1[0]; + b2 = smem2[0]; + + float2 delta; + delta.x = A12 * b2 - A22 * b1; + delta.y = A12 * b1 - A11 * b2; + + nextPt.x += delta.x; + nextPt.y += delta.y; + + if (::fabs(delta.x) < 0.01f && ::fabs(delta.y) < 0.01f) + break; + } + + float errval = 0; + if (calcErr) + { + for (int y = threadIdx.y, i = 0; y < c_multi_winSize_y::get(); y += blockDim.y, ++i) + { + for (int x = threadIdx.x, j = 0; x < c_multi_winSize_x::get(); x += blockDim.x, ++j) + { + work_type I_val = I_patch[i][j]; + work_type J_val = Tex_J_multi::read(nextPt.x + x + 0.5f, nextPt.y + y + 0.5f); + + work_type diff = J_val - I_val; + + accum(errval, abs_(diff)); + } + } + + reduce(smem1, errval, tid, plus()); + } + + if (tid == 0) + { + nextPt.x += c_multi_halfWin_x::get(); + nextPt.y += c_multi_halfWin_y::get(); + + nextPts[blockIdx.x] = nextPt; + + if (calcErr) + err[blockIdx.x] = static_cast(errval) / (cn * c_multi_winSize_x::get() * c_multi_winSize_y::get()); + } + } +#endif // defined(HAVE_TBB) + template void sparse_caller(int rows, int cols, const float2* prevPts, float2* nextPts, uchar* status, float* err, int ptcount, int level, dim3 block, cudaStream_t stream) @@ -326,6 +615,26 @@ namespace pyrlk cudaSafeCall( cudaDeviceSynchronize() ); } +#if defined(HAVE_TBB) + template + void sparse_caller_multi(int rows, int cols, const float2* prevPts, float2* nextPts, uchar* status, float* err, int ptcount, + int level, dim3 block, cudaStream_t stream) + { + dim3 grid(ptcount); + + if (level == 0 && err) + sparseKernel_multi<<>>(prevPts, nextPts, status, err, level, rows, cols); + else + sparseKernel_multi<<>>(prevPts, nextPts, status, err, level, rows, cols); + + cudaSafeCall( cudaGetLastError() ); + + if (stream == 0) + cudaSafeCall( cudaDeviceSynchronize() ); + } + +#endif // defined(HAVE_TBB) + template __global__ void denseKernel(PtrStepf u, PtrStepf v, const PtrStepf prevU, const PtrStepf prevV, PtrStepf err, const int rows, const int cols) { @@ -484,6 +793,30 @@ namespace pyrlk cudaSafeCall( cudaMemcpyToSymbol(c_iters, &iters, sizeof(int)) ); } +#if defined(HAVE_TBB) + void loadConstants_multi(int2 winSize, int iters, int index, cudaStream_t stream = 0) + { + int2 halfWin; +#define COPY_TO_SYMBOL_CALL(index) \ + cudaSafeCall( cudaMemcpyToSymbolAsync(c_winSize_x##index, &winSize.x, sizeof(int), 0, cudaMemcpyHostToDevice, stream) ); \ + cudaSafeCall( cudaMemcpyToSymbolAsync(c_winSize_y##index, &winSize.y, sizeof(int), 0, cudaMemcpyHostToDevice, stream) ); \ + halfWin = make_int2((winSize.x - 1) / 2, (winSize.y - 1) / 2); \ + cudaSafeCall( cudaMemcpyToSymbolAsync(c_halfWin_x##index, &halfWin.x, sizeof(int), 0, cudaMemcpyHostToDevice, stream) ); \ + cudaSafeCall( cudaMemcpyToSymbolAsync(c_halfWin_y##index, &halfWin.y, sizeof(int), 0, cudaMemcpyHostToDevice, stream) ); \ + cudaSafeCall( cudaMemcpyToSymbolAsync(c_iters##index, &iters, sizeof(int), 0, cudaMemcpyHostToDevice, stream) ); + + switch(index) + { + case 0: COPY_TO_SYMBOL_CALL(0) break; + case 1: COPY_TO_SYMBOL_CALL(1) break; + case 2: COPY_TO_SYMBOL_CALL(2) break; + case 3: COPY_TO_SYMBOL_CALL(3) break; + case 4: COPY_TO_SYMBOL_CALL(4) break; + default: CV_Error(CV_StsBadArg, "invalid execution line index"); break; + } + } +#endif // defined(HAVE_TBB) + void sparse1(PtrStepSzf I, PtrStepSzf J, const float2* prevPts, float2* nextPts, uchar* status, float* err, int ptcount, int level, dim3 block, dim3 patch, cudaStream_t stream) { @@ -528,6 +861,161 @@ namespace pyrlk level, block, stream); } +#if defined(HAVE_TBB) + void sparse1_multi(PtrStepSzf I, PtrStepSzf J, const float2* prevPts, float2* nextPts, uchar* status, float* err, int ptcount, + int level, dim3 block, dim3 patch, cudaStream_t stream, int index) + { + typedef void (*func_t)(int rows, int cols, const float2* prevPts, float2* nextPts, uchar* status, float* err, int ptcount, + int level, dim3 block, cudaStream_t stream); + + static const func_t funcs[5][5][5] = + { + { // index 0 + {sparse_caller_multi<1, 0, 1, 1>, sparse_caller_multi<1, 0, 2, 1>, sparse_caller_multi<1, 0, 3, 1>, sparse_caller_multi<1, 0, 4, 1>, sparse_caller_multi<1, 0, 5, 1>}, + {sparse_caller_multi<1, 0, 1, 2>, sparse_caller_multi<1, 0, 2, 2>, sparse_caller_multi<1, 0, 3, 2>, sparse_caller_multi<1, 0, 4, 2>, sparse_caller_multi<1, 0, 5, 2>}, + {sparse_caller_multi<1, 0, 1, 3>, sparse_caller_multi<1, 0, 2, 3>, sparse_caller_multi<1, 0, 3, 3>, sparse_caller_multi<1, 0, 4, 3>, sparse_caller_multi<1, 0, 5, 3>}, + {sparse_caller_multi<1, 0, 1, 4>, sparse_caller_multi<1, 0, 2, 4>, sparse_caller_multi<1, 0, 3, 4>, sparse_caller_multi<1, 0, 4, 4>, sparse_caller_multi<1, 0, 5, 4>}, + {sparse_caller_multi<1, 0, 1, 5>, sparse_caller_multi<1, 0, 2, 5>, sparse_caller_multi<1, 0, 3, 5>, sparse_caller_multi<1, 0, 4, 5>, sparse_caller_multi<1, 0, 5, 5>} + }, + { // index 1 + {sparse_caller_multi<1, 1, 1, 1>, sparse_caller_multi<1, 1, 2, 1>, sparse_caller_multi<1, 1, 3, 1>, sparse_caller_multi<1, 1, 4, 1>, sparse_caller_multi<1, 1, 5, 1>}, + {sparse_caller_multi<1, 1, 1, 2>, sparse_caller_multi<1, 1, 2, 2>, sparse_caller_multi<1, 1, 3, 2>, sparse_caller_multi<1, 1, 4, 2>, sparse_caller_multi<1, 1, 5, 2>}, + {sparse_caller_multi<1, 1, 1, 3>, sparse_caller_multi<1, 1, 2, 3>, sparse_caller_multi<1, 1, 3, 3>, sparse_caller_multi<1, 1, 4, 3>, sparse_caller_multi<1, 1, 5, 3>}, + {sparse_caller_multi<1, 1, 1, 4>, sparse_caller_multi<1, 1, 2, 4>, sparse_caller_multi<1, 1, 3, 4>, sparse_caller_multi<1, 1, 4, 4>, sparse_caller_multi<1, 1, 5, 4>}, + {sparse_caller_multi<1, 1, 1, 5>, sparse_caller_multi<1, 1, 2, 5>, sparse_caller_multi<1, 1, 3, 5>, sparse_caller_multi<1, 1, 4, 5>, sparse_caller_multi<1, 1, 5, 5>} + }, + { // index 2 + {sparse_caller_multi<1, 2, 1, 1>, sparse_caller_multi<1, 2, 2, 1>, sparse_caller_multi<1, 2, 3, 1>, sparse_caller_multi<1, 2, 4, 1>, sparse_caller_multi<1, 2, 5, 1>}, + {sparse_caller_multi<1, 2, 1, 2>, sparse_caller_multi<1, 2, 2, 2>, sparse_caller_multi<1, 2, 3, 2>, sparse_caller_multi<1, 2, 4, 2>, sparse_caller_multi<1, 2, 5, 2>}, + {sparse_caller_multi<1, 2, 1, 3>, sparse_caller_multi<1, 2, 2, 3>, sparse_caller_multi<1, 2, 3, 3>, sparse_caller_multi<1, 2, 4, 3>, sparse_caller_multi<1, 2, 5, 3>}, + {sparse_caller_multi<1, 2, 1, 4>, sparse_caller_multi<1, 2, 2, 4>, sparse_caller_multi<1, 2, 3, 4>, sparse_caller_multi<1, 2, 4, 4>, sparse_caller_multi<1, 2, 5, 4>}, + {sparse_caller_multi<1, 2, 1, 5>, sparse_caller_multi<1, 2, 2, 5>, sparse_caller_multi<1, 2, 3, 5>, sparse_caller_multi<1, 2, 4, 5>, sparse_caller_multi<1, 2, 5, 5>} + }, + { // index 3 + {sparse_caller_multi<1, 3, 1, 1>, sparse_caller_multi<1, 3, 2, 1>, sparse_caller_multi<1, 3, 3, 1>, sparse_caller_multi<1, 3, 4, 1>, sparse_caller_multi<1, 3, 5, 1>}, + {sparse_caller_multi<1, 3, 1, 2>, sparse_caller_multi<1, 3, 2, 2>, sparse_caller_multi<1, 3, 3, 2>, sparse_caller_multi<1, 3, 4, 2>, sparse_caller_multi<1, 3, 5, 2>}, + {sparse_caller_multi<1, 3, 1, 3>, sparse_caller_multi<1, 3, 2, 3>, sparse_caller_multi<1, 3, 3, 3>, sparse_caller_multi<1, 3, 4, 3>, sparse_caller_multi<1, 3, 5, 3>}, + {sparse_caller_multi<1, 3, 1, 4>, sparse_caller_multi<1, 3, 2, 4>, sparse_caller_multi<1, 3, 3, 4>, sparse_caller_multi<1, 3, 4, 4>, sparse_caller_multi<1, 3, 5, 4>}, + {sparse_caller_multi<1, 3, 1, 5>, sparse_caller_multi<1, 3, 2, 5>, sparse_caller_multi<1, 3, 3, 5>, sparse_caller_multi<1, 3, 4, 5>, sparse_caller_multi<1, 3, 5, 5>} + }, + { // index 4 + {sparse_caller_multi<1, 4, 1, 1>, sparse_caller_multi<1, 4, 2, 1>, sparse_caller_multi<1, 4, 3, 1>, sparse_caller_multi<1, 4, 4, 1>, sparse_caller_multi<1, 4, 5, 1>}, + {sparse_caller_multi<1, 4, 1, 2>, sparse_caller_multi<1, 4, 2, 2>, sparse_caller_multi<1, 4, 3, 2>, sparse_caller_multi<1, 4, 4, 2>, sparse_caller_multi<1, 4, 5, 2>}, + {sparse_caller_multi<1, 4, 1, 3>, sparse_caller_multi<1, 4, 2, 3>, sparse_caller_multi<1, 4, 3, 3>, sparse_caller_multi<1, 4, 4, 3>, sparse_caller_multi<1, 4, 5, 3>}, + {sparse_caller_multi<1, 4, 1, 4>, sparse_caller_multi<1, 4, 2, 4>, sparse_caller_multi<1, 4, 3, 4>, sparse_caller_multi<1, 4, 4, 4>, sparse_caller_multi<1, 4, 5, 4>}, + {sparse_caller_multi<1, 4, 1, 5>, sparse_caller_multi<1, 4, 2, 5>, sparse_caller_multi<1, 4, 3, 5>, sparse_caller_multi<1, 4, 4, 5>, sparse_caller_multi<1, 4, 5, 5>} + } + }; + + switch(index) + { + case 0: + bindTexture(&tex_If1_multi0, I); + bindTexture(&tex_Jf1_multi0, J); + break; + case 1: + bindTexture(&tex_If1_multi1, I); + bindTexture(&tex_Jf1_multi1, J); + break; + case 2: + bindTexture(&tex_If1_multi2, I); + bindTexture(&tex_Jf1_multi2, J); + break; + case 3: + bindTexture(&tex_If1_multi3, I); + bindTexture(&tex_Jf1_multi3, J); + break; + case 4: + bindTexture(&tex_If1_multi4, I); + bindTexture(&tex_Jf1_multi4, J); + break; + default: + CV_Error(CV_StsBadArg, "invalid execution line index"); + break; + } + + funcs[index][patch.y - 1][patch.x - 1](I.rows, I.cols, prevPts, nextPts, status, err, ptcount, + level, block, stream); + } + + void sparse4_multi(PtrStepSz I, PtrStepSz J, const float2* prevPts, float2* nextPts, uchar* status, float* err, int ptcount, + int level, dim3 block, dim3 patch, cudaStream_t stream, int index) + { + typedef void (*func_t)(int rows, int cols, const float2* prevPts, float2* nextPts, uchar* status, float* err, int ptcount, + int level, dim3 block, cudaStream_t stream); + + static const func_t funcs[5][5][5] = + { + { // index 0 + {sparse_caller_multi<4, 0, 1, 1>, sparse_caller_multi<4, 0, 2, 1>, sparse_caller_multi<4, 0, 3, 1>, sparse_caller_multi<4, 0, 4, 1>, sparse_caller_multi<4, 0, 5, 1>}, + {sparse_caller_multi<4, 0, 1, 2>, sparse_caller_multi<4, 0, 2, 2>, sparse_caller_multi<4, 0, 3, 2>, sparse_caller_multi<4, 0, 4, 2>, sparse_caller_multi<4, 0, 5, 2>}, + {sparse_caller_multi<4, 0, 1, 3>, sparse_caller_multi<4, 0, 2, 3>, sparse_caller_multi<4, 0, 3, 3>, sparse_caller_multi<4, 0, 4, 3>, sparse_caller_multi<4, 0, 5, 3>}, + {sparse_caller_multi<4, 0, 1, 4>, sparse_caller_multi<4, 0, 2, 4>, sparse_caller_multi<4, 0, 3, 4>, sparse_caller_multi<4, 0, 4, 4>, sparse_caller_multi<4, 0, 5, 4>}, + {sparse_caller_multi<4, 0, 1, 5>, sparse_caller_multi<4, 0, 2, 5>, sparse_caller_multi<4, 0, 3, 5>, sparse_caller_multi<4, 0, 4, 5>, sparse_caller_multi<4, 0, 5, 5>} + }, + { // index 1 + {sparse_caller_multi<4, 1, 1, 1>, sparse_caller_multi<4, 1, 2, 1>, sparse_caller_multi<4, 1, 3, 1>, sparse_caller_multi<4, 1, 4, 1>, sparse_caller_multi<4, 1, 5, 1>}, + {sparse_caller_multi<4, 1, 1, 2>, sparse_caller_multi<4, 1, 2, 2>, sparse_caller_multi<4, 1, 3, 2>, sparse_caller_multi<4, 1, 4, 2>, sparse_caller_multi<4, 1, 5, 2>}, + {sparse_caller_multi<4, 1, 1, 3>, sparse_caller_multi<4, 1, 2, 3>, sparse_caller_multi<4, 1, 3, 3>, sparse_caller_multi<4, 1, 4, 3>, sparse_caller_multi<4, 1, 5, 3>}, + {sparse_caller_multi<4, 1, 1, 4>, sparse_caller_multi<4, 1, 2, 4>, sparse_caller_multi<4, 1, 3, 4>, sparse_caller_multi<4, 1, 4, 4>, sparse_caller_multi<4, 1, 5, 4>}, + {sparse_caller_multi<4, 1, 1, 5>, sparse_caller_multi<4, 1, 2, 5>, sparse_caller_multi<4, 1, 3, 5>, sparse_caller_multi<4, 1, 4, 5>, sparse_caller_multi<4, 1, 5, 5>} + }, + { // index 2 + {sparse_caller_multi<4, 2, 1, 1>, sparse_caller_multi<4, 2, 2, 1>, sparse_caller_multi<4, 2, 3, 1>, sparse_caller_multi<4, 2, 4, 1>, sparse_caller_multi<4, 2, 5, 1>}, + {sparse_caller_multi<4, 2, 1, 2>, sparse_caller_multi<4, 2, 2, 2>, sparse_caller_multi<4, 2, 3, 2>, sparse_caller_multi<4, 2, 4, 2>, sparse_caller_multi<4, 2, 5, 2>}, + {sparse_caller_multi<4, 2, 1, 3>, sparse_caller_multi<4, 2, 2, 3>, sparse_caller_multi<4, 2, 3, 3>, sparse_caller_multi<4, 2, 4, 3>, sparse_caller_multi<4, 2, 5, 3>}, + {sparse_caller_multi<4, 2, 1, 4>, sparse_caller_multi<4, 2, 2, 4>, sparse_caller_multi<4, 2, 3, 4>, sparse_caller_multi<4, 2, 4, 4>, sparse_caller_multi<4, 2, 5, 4>}, + {sparse_caller_multi<4, 2, 1, 5>, sparse_caller_multi<4, 2, 2, 5>, sparse_caller_multi<4, 2, 3, 5>, sparse_caller_multi<4, 2, 4, 5>, sparse_caller_multi<4, 2, 5, 5>} + }, + { // index 3 + {sparse_caller_multi<4, 3, 1, 1>, sparse_caller_multi<4, 3, 2, 1>, sparse_caller_multi<4, 3, 3, 1>, sparse_caller_multi<4, 3, 4, 1>, sparse_caller_multi<4, 3, 5, 1>}, + {sparse_caller_multi<4, 3, 1, 2>, sparse_caller_multi<4, 3, 2, 2>, sparse_caller_multi<4, 3, 3, 2>, sparse_caller_multi<4, 3, 4, 2>, sparse_caller_multi<4, 3, 5, 2>}, + {sparse_caller_multi<4, 3, 1, 3>, sparse_caller_multi<4, 3, 2, 3>, sparse_caller_multi<4, 3, 3, 3>, sparse_caller_multi<4, 3, 4, 3>, sparse_caller_multi<4, 3, 5, 3>}, + {sparse_caller_multi<4, 3, 1, 4>, sparse_caller_multi<4, 3, 2, 4>, sparse_caller_multi<4, 3, 3, 4>, sparse_caller_multi<4, 3, 4, 4>, sparse_caller_multi<4, 3, 5, 4>}, + {sparse_caller_multi<4, 3, 1, 5>, sparse_caller_multi<4, 3, 2, 5>, sparse_caller_multi<4, 3, 3, 5>, sparse_caller_multi<4, 3, 4, 5>, sparse_caller_multi<4, 3, 5, 5>} + }, + { // index 4 + {sparse_caller_multi<4, 4, 1, 1>, sparse_caller_multi<4, 4, 2, 1>, sparse_caller_multi<4, 4, 3, 1>, sparse_caller_multi<4, 4, 4, 1>, sparse_caller_multi<4, 4, 5, 1>}, + {sparse_caller_multi<4, 4, 1, 2>, sparse_caller_multi<4, 4, 2, 2>, sparse_caller_multi<4, 4, 3, 2>, sparse_caller_multi<4, 4, 4, 2>, sparse_caller_multi<4, 4, 5, 2>}, + {sparse_caller_multi<4, 4, 1, 3>, sparse_caller_multi<4, 4, 2, 3>, sparse_caller_multi<4, 4, 3, 3>, sparse_caller_multi<4, 4, 4, 3>, sparse_caller_multi<4, 4, 5, 3>}, + {sparse_caller_multi<4, 4, 1, 4>, sparse_caller_multi<4, 4, 2, 4>, sparse_caller_multi<4, 4, 3, 4>, sparse_caller_multi<4, 4, 4, 4>, sparse_caller_multi<4, 4, 5, 4>}, + {sparse_caller_multi<4, 4, 1, 5>, sparse_caller_multi<4, 4, 2, 5>, sparse_caller_multi<4, 4, 3, 5>, sparse_caller_multi<4, 4, 4, 5>, sparse_caller_multi<4, 4, 5, 5>} + } + }; + + switch(index) + { + case 0: + bindTexture(&tex_If4_multi0, I); + bindTexture(&tex_Jf4_multi0, J); + break; + case 1: + bindTexture(&tex_If4_multi1, I); + bindTexture(&tex_Jf4_multi1, J); + break; + case 2: + bindTexture(&tex_If4_multi2, I); + bindTexture(&tex_Jf4_multi2, J); + break; + case 3: + bindTexture(&tex_If4_multi3, I); + bindTexture(&tex_Jf4_multi3, J); + break; + case 4: + bindTexture(&tex_If4_multi4, I); + bindTexture(&tex_Jf4_multi4, J); + break; + default: + CV_Error(CV_StsBadArg, "invalid execution line index"); + break; + } + + funcs[index][patch.y - 1][patch.x - 1](I.rows, I.cols, prevPts, nextPts, status, err, ptcount, + level, block, stream); + } + +#endif // defined(HAVE_TBB) + void dense(PtrStepSzb I, PtrStepSzf J, PtrStepSzf u, PtrStepSzf v, PtrStepSzf prevU, PtrStepSzf prevV, PtrStepSzf err, int2 winSize, cudaStream_t stream) { dim3 block(16, 16); diff --git a/modules/gpu/src/pyrlk.cpp b/modules/gpu/src/pyrlk.cpp index b6c925a111..052db3b61e 100644 --- a/modules/gpu/src/pyrlk.cpp +++ b/modules/gpu/src/pyrlk.cpp @@ -42,6 +42,11 @@ #include "precomp.hpp" +#ifdef HAVE_TBB +#include +#include +#endif + using namespace std; using namespace cv; using namespace cv::gpu; @@ -64,6 +69,22 @@ namespace pyrlk void sparse4(PtrStepSz I, PtrStepSz J, const float2* prevPts, float2* nextPts, uchar* status, float* err, int ptcount, int level, dim3 block, dim3 patch, cudaStream_t stream = 0); + +#if !defined(HAVE_TBB) + void loadConstants_multi(int2, int, int, cudaStream_t) { throw_notbb(); } + void sparse1_multi(PtrStepSzf, PtrStepSzf, const float2*, float2*, uchar*, float*, int, + int, dim3, dim3, cudaStream_t, int) { throw_notbb(); } + void sparse4_multi(PtrStepSz, PtrStepSz, const float2*, float2*, uchar*, float*, int, + int, dim3, dim3, cudaStream_t, int) { throw_notbb(); } +#else + void loadConstants_multi(int2 winSize, int iters, int index = 0, cudaStream_t stream = 0); + + void sparse1_multi(PtrStepSzf I, PtrStepSzf J, const float2* prevPts, float2* nextPts, uchar* status, float* err, int ptcount, + int level, dim3 block, dim3 patch, cudaStream_t stream = 0, int index = 0); + void sparse4_multi(PtrStepSz I, PtrStepSz J, const float2* prevPts, float2* nextPts, uchar* status, float* err, int ptcount, + int level, dim3 block, dim3 patch, cudaStream_t stream = 0, int index = 0); +#endif + void dense(PtrStepSzb I, PtrStepSzf J, PtrStepSzf u, PtrStepSzf v, PtrStepSzf prevU, PtrStepSzf prevV, PtrStepSzf err, int2 winSize, cudaStream_t stream = 0); } @@ -98,7 +119,9 @@ namespace } } -void cv::gpu::PyrLKOpticalFlow::sparse(const GpuMat& prevImg, const GpuMat& nextImg, const GpuMat& prevPts, GpuMat& nextPts, GpuMat& status, GpuMat* err) +void cv::gpu::PyrLKOpticalFlow::sparse(const GpuMat& prevImg, + const GpuMat& nextImg, const GpuMat& prevPts, + GpuMat& nextPts, GpuMat& status, GpuMat* err) { if (prevPts.empty()) { @@ -181,6 +204,130 @@ void cv::gpu::PyrLKOpticalFlow::sparse(const GpuMat& prevImg, const GpuMat& next } } +#ifdef HAVE_TBB +//-------------------------------------------------------------------------- +// Multi-threading support + +static bool index_vector_use[5] = {true, true, true, true, true}; // all free +static tbb::mutex s_PyrLKOpticalFlow_Mutex; +static condition_variable s_PyrLKOpticalFlow_ConditionVariable; + +void cv::gpu::PyrLKOpticalFlow::sparse_multi(const GpuMat& prevImg, + const GpuMat& nextImg, const GpuMat& prevPts, GpuMat& nextPts, + GpuMat& status, Stream& stream, GpuMat* err) +{ + if (prevPts.empty()) + { + nextPts.release(); + status.release(); + if (err) err->release(); + return; + } + + dim3 block, patch; + calcPatchSize(winSize, block, patch); + + CV_Assert(prevImg.channels() == 1 || prevImg.channels() == 3 || prevImg.channels() == 4); + CV_Assert(prevImg.size() == nextImg.size() && prevImg.type() == nextImg.type()); + CV_Assert(maxLevel >= 0); + CV_Assert(winSize.width > 2 && winSize.height > 2); + CV_Assert(patch.x > 0 && patch.x < 6 && patch.y > 0 && patch.y < 6); + CV_Assert(prevPts.rows == 1 && prevPts.type() == CV_32FC2); + + if (useInitialFlow) + CV_Assert(nextPts.size() == prevPts.size() && nextPts.type() == CV_32FC2); + else + ensureSizeIsEnough(1, prevPts.cols, prevPts.type(), nextPts); + + GpuMat temp1 = (useInitialFlow ? nextPts : prevPts).reshape(1); + GpuMat temp2 = nextPts.reshape(1); + multiply(temp1, Scalar::all(1.0 / (1 << maxLevel) / 2.0), temp2); + + ensureSizeIsEnough(1, prevPts.cols, CV_8UC1, status); + status.setTo(Scalar::all(1)); + + if (err) + ensureSizeIsEnough(1, prevPts.cols, CV_32FC1, *err); + + // build the image pyramids. + + prevPyr_.resize(maxLevel + 1); + nextPyr_.resize(maxLevel + 1); + + int cn = prevImg.channels(); + + if (cn == 1 || cn == 4) + { + prevImg.convertTo(prevPyr_[0], CV_32F); + nextImg.convertTo(nextPyr_[0], CV_32F); + } + else + { + buf_.resize(1); + + cvtColor(prevImg, buf_[0], COLOR_BGR2BGRA); + buf_[0].convertTo(prevPyr_[0], CV_32F); + + cvtColor(nextImg, buf_[0], COLOR_BGR2BGRA); + buf_[0].convertTo(nextPyr_[0], CV_32F); + } + + for (int level = 1; level <= maxLevel; ++level) + { + pyrDown(prevPyr_[level - 1], prevPyr_[level]); + pyrDown(nextPyr_[level - 1], nextPyr_[level]); + } + + //-------------------------------------------------------------------------- + // Multithreading support + int index = -1; + + do + { + unique_lock ul(s_PyrLKOpticalFlow_Mutex); + for (unsigned int uiI = 0; uiI < 5; ++uiI) + { + if (index_vector_use[uiI]) + { + index = uiI; + index_vector_use[uiI] = false; + break; + } + } + if (index < 0) + s_PyrLKOpticalFlow_ConditionVariable.wait(ul); + + ul.unlock(); + }while (index < 0); + + //-------------------------------------------------------------------------- + + pyrlk::loadConstants_multi(make_int2(winSize.width, winSize.height), iters, index); + + for (int level = maxLevel; level >= 0; level--) + { + if (cn == 1) + { + pyrlk::sparse1_multi(prevPyr_[level], nextPyr_[level], + prevPts.ptr(), nextPts.ptr(), status.ptr(), + level == 0 && err ? err->ptr() : 0, prevPts.cols, + level, block, patch, StreamAccessor::getStream(stream), index); + } + else + { + pyrlk::sparse4_multi(prevPyr_[level], nextPyr_[level], + prevPts.ptr(), nextPts.ptr(), status.ptr(), + level == 0 && err ? err->ptr() : 0, prevPts.cols, + level, block, patch, StreamAccessor::getStream(stream), index); + } + } + + unique_lock ul(s_PyrLKOpticalFlow_Mutex); + index_vector_use[index] = true; + s_PyrLKOpticalFlow_ConditionVariable.notify_one(); +} +#endif + void cv::gpu::PyrLKOpticalFlow::dense(const GpuMat& prevImg, const GpuMat& nextImg, GpuMat& u, GpuMat& v, GpuMat* err) { CV_Assert(prevImg.type() == CV_8UC1); diff --git a/modules/gpu/test/test_optflow.cpp b/modules/gpu/test/test_optflow.cpp index 343a756b3f..38b166d35a 100644 --- a/modules/gpu/test/test_optflow.cpp +++ b/modules/gpu/test/test_optflow.cpp @@ -44,6 +44,10 @@ #ifdef HAVE_CUDA +#ifdef HAVE_TBB +#include +#endif + using namespace cvtest; ////////////////////////////////////////////////////// @@ -322,6 +326,134 @@ GPU_TEST_P(PyrLKOpticalFlow, Sparse) ASSERT_LE(bad_ratio, 0.01); } +#ifdef HAVE_TBB + +struct Sparse_Multi_Functor +{ + explicit Sparse_Multi_Functor(const cv::Mat& in_frame0, const cv::Mat& in_frame1, + const cv::Mat& in_pts_mat, + cv::gpu::GpuMat* in_d_pts, + cv::gpu::GpuMat* in_d_nextPts, + cv::gpu::GpuMat* in_d_status, + cv::gpu::Stream* in_streams): + m_frame0(in_frame0), m_frame1(in_frame1), + m_pts_mat(in_pts_mat), + m_d_pts(in_d_pts), m_d_nextPts(in_d_nextPts), + m_d_status(in_d_status), m_streams(in_streams){} + + void operator()( const tbb::blocked_range& r ) const + { + for( size_t i = r.begin(); i != r.end(); ++i ) + { + m_d_pts[i].upload(m_pts_mat); + cv::gpu::PyrLKOpticalFlow pyrLK; + pyrLK.sparse_multi(loadMat(m_frame0), loadMat(m_frame1), m_d_pts[i], + m_d_nextPts[i], m_d_status[i], m_streams[i]); + m_streams[i].waitForCompletion(); + } + } + + const cv::Mat& m_frame0; + const cv::Mat& m_frame1; + const cv::Mat& m_pts_mat; + cv::gpu::GpuMat* m_d_pts; + cv::gpu::GpuMat* m_d_nextPts; + cv::gpu::GpuMat* m_d_status; + cv::gpu::Stream* m_streams; +}; + + +GPU_TEST_P(PyrLKOpticalFlow, Sparse_Multi) +{ + cv::Mat frame0 = readImage("opticalflow/frame0.png", useGray ? cv::IMREAD_GRAYSCALE : cv::IMREAD_COLOR); + ASSERT_FALSE(frame0.empty()); + + cv::Mat frame1 = readImage("opticalflow/frame1.png", useGray ? cv::IMREAD_GRAYSCALE : cv::IMREAD_COLOR); + ASSERT_FALSE(frame1.empty()); + + cv::Mat gray_frame; + if (useGray) + gray_frame = frame0; + else + cv::cvtColor(frame0, gray_frame, cv::COLOR_BGR2GRAY); + + std::vector pts; + cv::goodFeaturesToTrack(gray_frame, pts, 1000, 0.01, 0.0); + + //-------------------------------------------------------------------------- + // GPU + const unsigned int NB_EXEC_LINES = 27; + + cv::gpu::GpuMat d_pts[NB_EXEC_LINES]; + cv::gpu::GpuMat d_nextPts[NB_EXEC_LINES]; + cv::gpu::GpuMat d_status[NB_EXEC_LINES]; + cv::gpu::Stream streams[NB_EXEC_LINES]; + + cv::Mat pts_mat(1, (int) pts.size(), CV_32FC2, (void*) &pts[0]); + + tbb::parallel_for(tbb::blocked_range(0, NB_EXEC_LINES), + Sparse_Multi_Functor(frame0, frame1, pts_mat, + d_pts, d_nextPts, d_status, streams)); + + std::vector nextPts[NB_EXEC_LINES]; + std::vector status[NB_EXEC_LINES]; + + for (unsigned int i = 0; i < NB_EXEC_LINES; ++i) + { + nextPts[i].resize(d_nextPts[i].cols); + cv::Mat nextPts_mat(1, d_nextPts[i].cols, CV_32FC2, (void*) &(nextPts[i][0])); + d_nextPts[i].download(nextPts_mat); + + status[i].resize(d_status[i].cols); + cv::Mat status_mat(1, d_status[i].cols, CV_8UC1, (void*) &(status[i][0])); + d_status[i].download(status_mat); + } + + //-------------------------------------------------------------------------- + // CPU + std::vector nextPts_gold; + std::vector status_gold; + cv::calcOpticalFlowPyrLK(frame0, frame1, pts, nextPts_gold, status_gold, cv::noArray()); + + //-------------------------------------------------------------------------- + // CHECKS + for (unsigned int uiI = 0; uiI < NB_EXEC_LINES; ++uiI) + { + ASSERT_EQ(nextPts_gold.size(), nextPts[uiI].size()); + ASSERT_EQ(status_gold.size(), status[uiI].size()); + } + + size_t mistmatch = 0; + for (unsigned int uiI = 0; uiI < NB_EXEC_LINES; ++uiI) + { + for (size_t i = 0; i < nextPts[uiI].size(); ++i) + { + cv::Point2i a = nextPts[uiI][i]; + cv::Point2i b = nextPts_gold[i]; + + if (status[uiI][i] != status_gold[i]) + { + ++mistmatch; + continue; + } + + if (status[uiI][i]) + { + bool eq = std::abs(a.x - b.x) <= 1 && std::abs(a.y - b.y) <= 1; + + if (!eq) + ++mistmatch; + } + } + } + + double bad_ratio = static_cast(mistmatch) / (nextPts[0].size() * NB_EXEC_LINES); + + ASSERT_LE(bad_ratio, 0.01); +} + +#endif // HAVE_TBB + INSTANTIATE_TEST_CASE_P(GPU_Video, PyrLKOpticalFlow, testing::Combine( ALL_DEVICES, testing::Values(UseGray(true), UseGray(false))));