diff --git a/modules/cudaoptflow/perf/perf_optflow.cpp b/modules/cudaoptflow/perf/perf_optflow.cpp index 57994b7f4..d2992c30c 100644 --- a/modules/cudaoptflow/perf/perf_optflow.cpp +++ b/modules/cudaoptflow/perf/perf_optflow.cpp @@ -116,10 +116,10 @@ PERF_TEST_P(ImagePair_Gray_NPts_WinSz_Levels_Iters, PyrLKOpticalFlowSparse, 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); + 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); + cv::Mat frame1 = readImage(imagePair.second, useGray ? cv::IMREAD_GRAYSCALE : cv::IMREAD_COLOR); ASSERT_FALSE(frame1.empty()); cv::Mat gray_frame; @@ -131,6 +131,14 @@ PERF_TEST_P(ImagePair_Gray_NPts_WinSz_Levels_Iters, PyrLKOpticalFlowSparse, cv::Mat pts; cv::goodFeaturesToTrack(gray_frame, pts, points, 0.01, 0.0); + frame0.convertTo(frame0, CV_32F); + frame1.convertTo(frame1, CV_32F); + if(!useGray) + { + cv::cvtColor(frame0, frame0, cv::COLOR_BGR2BGRA); + cv::cvtColor(frame1, frame1, cv::COLOR_BGR2BGRA); + } + if (PERF_RUN_CUDA()) { const cv::cuda::GpuMat d_pts(pts.reshape(2, 1)); @@ -318,4 +326,4 @@ PERF_TEST_P(ImagePair, OpticalFlowDual_TVL1, CPU_SANITY_CHECK(flow); } -} +} \ No newline at end of file diff --git a/modules/cudaoptflow/src/cuda/pyrlk.cu b/modules/cudaoptflow/src/cuda/pyrlk.cu index 7693551fc..5d40a47ea 100644 --- a/modules/cudaoptflow/src/cuda/pyrlk.cu +++ b/modules/cudaoptflow/src/cuda/pyrlk.cu @@ -48,6 +48,8 @@ #include "opencv2/core/cuda/limits.hpp" #include "opencv2/core/cuda/vec_math.hpp" #include "opencv2/core/cuda/reduce.hpp" +#include "opencv2/core/cuda/filters.hpp" +#include "opencv2/core/cuda/border_interpolate.hpp" using namespace cv::cuda; using namespace cv::cuda::device; @@ -60,53 +62,240 @@ namespace pyrlk __constant__ int c_halfWin_y; __constant__ int c_iters; + texture tex_I8U(false, cudaFilterModeLinear, cudaAddressModeClamp); + texture tex_I8UC4(false, cudaFilterModeLinear, cudaAddressModeClamp); + + texture tex_I16UC4(false, cudaFilterModeLinear, cudaAddressModeClamp); + + texture tex_If(false, cudaFilterModeLinear, cudaAddressModeClamp); texture tex_If4(false, cudaFilterModeLinear, cudaAddressModeClamp); + texture tex_Ib(false, cudaFilterModePoint, cudaAddressModeClamp); + texture tex_J8U(false, cudaFilterModeLinear, cudaAddressModeClamp); + texture tex_J8UC4(false, cudaFilterModeLinear, cudaAddressModeClamp); + + texture tex_J16UC4(false, cudaFilterModeLinear, cudaAddressModeClamp); + + texture tex_Jf(false, cudaFilterModeLinear, cudaAddressModeClamp); texture tex_Jf4(false, cudaFilterModeLinear, cudaAddressModeClamp); - template struct Tex_I; - template <> struct Tex_I<1> + + template struct Tex_I + { + static __host__ __forceinline__ void bindTexture_(PtrStepSz::vec_type> I) + { + (void)I; + } + }; + + template <> struct Tex_I<1, uchar> + { + static __device__ __forceinline__ float read(float x, float y) + { + return tex2D(tex_I8U, x, y); + } + static __host__ __forceinline__ void bindTexture_(PtrStepSz& I) + { + bindTexture(&tex_I8U, I); + } + }; + template <> struct Tex_I<1, ushort> + { + static __device__ __forceinline__ float read(float x, float y) + { + return 0.0; + } + static __host__ __forceinline__ void bindTexture_(PtrStepSz& I) + { + (void)I; + } + }; + template <> struct Tex_I<1, int> + { + static __device__ __forceinline__ float read(float x, float y) + { + return 0.0; + } + static __host__ __forceinline__ void bindTexture_(PtrStepSz& I) + { + (void)I; + } + }; + template <> struct Tex_I<1, float> { static __device__ __forceinline__ float read(float x, float y) { return tex2D(tex_If, x, y); } + static __host__ __forceinline__ void bindTexture_(PtrStepSz& I) + { + bindTexture(&tex_If, I); + } }; - template <> struct Tex_I<4> + // ****************** 3 channel specializations ************************ + template <> struct Tex_I<3, uchar> + { + static __device__ __forceinline__ float3 read(float x, float y) + { + return make_float3(0,0,0); + } + static __host__ __forceinline__ void bindTexture_(PtrStepSz I) + { + (void)I; + } + }; + template <> struct Tex_I<3, ushort> + { + static __device__ __forceinline__ float3 read(float x, float y) + { + return make_float3(0, 0, 0); + } + static __host__ __forceinline__ void bindTexture_(PtrStepSz I) + { + (void)I; + } + }; + template <> struct Tex_I<3, int> + { + static __device__ __forceinline__ float3 read(float x, float y) + { + return make_float3(0, 0, 0); + } + static __host__ __forceinline__ void bindTexture_(PtrStepSz I) + { + (void)I; + } + }; + template <> struct Tex_I<3, float> + { + static __device__ __forceinline__ float3 read(float x, float y) + { + return make_float3(0, 0, 0); + } + static __host__ __forceinline__ void bindTexture_(PtrStepSz I) + { + (void)I; + } + }; + // ****************** 4 channel specializations ************************ + + template <> struct Tex_I<4, uchar> + { + static __device__ __forceinline__ float4 read(float x, float y) + { + return tex2D(tex_I8UC4, x, y); + } + static __host__ __forceinline__ void bindTexture_(PtrStepSz& I) + { + bindTexture(&tex_I8UC4, I); + } + }; + template <> struct Tex_I<4, ushort> + { + static __device__ __forceinline__ float4 read(float x, float y) + { + return tex2D(tex_I16UC4, x, y); + } + static __host__ __forceinline__ void bindTexture_(PtrStepSz& I) + { + bindTexture(&tex_I16UC4, I); + } + }; + template <> struct Tex_I<4, float> { static __device__ __forceinline__ float4 read(float x, float y) { return tex2D(tex_If4, x, y); } + static __host__ __forceinline__ void bindTexture_(PtrStepSz& I) + { + bindTexture(&tex_If4, I); + } }; - - template struct Tex_J; - template <> struct Tex_J<1> + // ************* J *************** + template struct Tex_J + { + static __host__ __forceinline__ void bindTexture_(PtrStepSz::vec_type>& J) + { + (void)J; + } + }; + template <> struct Tex_J<1, uchar> + { + static __device__ __forceinline__ float read(float x, float y) + { + return tex2D(tex_J8U, x, y); + } + static __host__ __forceinline__ void bindTexture_(PtrStepSz& J) + { + bindTexture(&tex_J8U, J); + } + }; + template <> struct Tex_J<1, float> { static __device__ __forceinline__ float read(float x, float y) { return tex2D(tex_Jf, x, y); } + static __host__ __forceinline__ void bindTexture_(PtrStepSz& J) + { + bindTexture(&tex_Jf, J); + } }; - template <> struct Tex_J<4> + // ************* 4 channel specializations *************** + template <> struct Tex_J<4, uchar> + { + static __device__ __forceinline__ float4 read(float x, float y) + { + return tex2D(tex_J8UC4, x, y); + } + static __host__ __forceinline__ void bindTexture_(PtrStepSz& J) + { + bindTexture(&tex_J8UC4, J); + } + }; + template <> struct Tex_J<4, ushort> + { + static __device__ __forceinline__ float4 read(float x, float y) + { + return tex2D(tex_J16UC4, x, y); + } + static __host__ __forceinline__ void bindTexture_(PtrStepSz& J) + { + bindTexture(&tex_J16UC4, J); + } + }; + template <> struct Tex_J<4, float> { static __device__ __forceinline__ float4 read(float x, float y) { return tex2D(tex_Jf4, x, y); } + static __host__ __forceinline__ void bindTexture_(PtrStepSz& J) + { + bindTexture(&tex_Jf4, J); + } }; - __device__ __forceinline__ void accum(float& dst, float val) + __device__ __forceinline__ void accum(float& dst, const float& val) { dst += val; } - __device__ __forceinline__ void accum(float& dst, const float4& val) + __device__ __forceinline__ void accum(float& dst, const float2& val) + { + dst += val.x + val.y; + } + __device__ __forceinline__ void accum(float& dst, const float3& val) { dst += val.x + val.y + val.z; } + __device__ __forceinline__ void accum(float& dst, const float4& val) + { + dst += val.x + val.y + val.z + val.w; + } __device__ __forceinline__ float abs_(float a) { @@ -116,8 +305,46 @@ namespace pyrlk { return abs(a); } + __device__ __forceinline__ float2 abs_(const float2& a) + { + return abs(a); + } + __device__ __forceinline__ float3 abs_(const float3& a) + { + return abs(a); + } - template + + template __device__ __forceinline__ typename TypeVec::vec_type ToFloat(const typename TypeVec::vec_type& other) + { + return other; + } + template __device__ __forceinline__ typename TypeVec::vec_type ToFloat(const typename TypeVec::vec_type& other) + { + typename TypeVec::vec_type ret; + ret.x = other.x; + ret.y = other.y; + return ret; + } + template __device__ __forceinline__ typename TypeVec::vec_type ToFloat(const typename TypeVec::vec_type& other) + { + typename TypeVec::vec_type ret; + ret.x = other.x; + ret.y = other.y; + ret.z = other.z; + return ret; + } + template __device__ __forceinline__ typename TypeVec::vec_type ToFloat(const typename TypeVec::vec_type& other) + { + typename TypeVec::vec_type ret; + ret.x = other.x; + ret.y = other.y; + ret.z = other.z; + ret.w = other.w; + return ret; + } + + template __global__ void sparseKernel(const float2* prevPts, float2* nextPts, uchar* status, float* err, const int level, const int rows, const int cols) { #if __CUDA_ARCH__ <= 110 @@ -166,15 +393,15 @@ namespace pyrlk float x = prevPt.x + xBase + 0.5f; float y = prevPt.y + yBase + 0.5f; - I_patch[i][j] = Tex_I::read(x, y); + I_patch[i][j] = Tex_I::read(x, y); // Sharr Deriv - work_type dIdx = 3.0f * Tex_I::read(x+1, y-1) + 10.0f * Tex_I::read(x+1, y) + 3.0f * Tex_I::read(x+1, y+1) - - (3.0f * Tex_I::read(x-1, y-1) + 10.0f * Tex_I::read(x-1, y) + 3.0f * Tex_I::read(x-1, y+1)); + work_type dIdx = 3.0f * Tex_I::read(x+1, y-1) + 10.0f * Tex_I::read(x+1, y) + 3.0f * Tex_I::read(x+1, y+1) - + (3.0f * Tex_I::read(x-1, y-1) + 10.0f * Tex_I::read(x-1, y) + 3.0f * Tex_I::read(x-1, y+1)); - work_type dIdy = 3.0f * Tex_I::read(x-1, y+1) + 10.0f * Tex_I::read(x, y+1) + 3.0f * Tex_I::read(x+1, y+1) - - (3.0f * Tex_I::read(x-1, y-1) + 10.0f * Tex_I::read(x, y-1) + 3.0f * Tex_I::read(x+1, y-1)); + work_type dIdy = 3.0f * Tex_I::read(x-1, y+1) + 10.0f * Tex_I::read(x, y+1) + 3.0f * Tex_I::read(x+1, y+1) - + (3.0f * Tex_I::read(x-1, y-1) + 10.0f * Tex_I::read(x, y-1) + 3.0f * Tex_I::read(x+1, y-1)); dIdx_patch[i][j] = dIdx; dIdy_patch[i][j] = dIdy; @@ -243,7 +470,7 @@ namespace pyrlk for (int x = threadIdx.x, j = 0; x < c_winSize_x; x += blockDim.x, ++j) { work_type I_val = I_patch[i][j]; - work_type J_val = Tex_J::read(nextPt.x + x + 0.5f, nextPt.y + y + 0.5f); + work_type J_val = Tex_J::read(nextPt.x + x + 0.5f, nextPt.y + y + 0.5f); work_type diff = (J_val - I_val) * 32.0f; @@ -286,7 +513,7 @@ namespace pyrlk for (int x = threadIdx.x, j = 0; x < c_winSize_x; x += blockDim.x, ++j) { work_type I_val = I_patch[i][j]; - work_type J_val = Tex_J::read(nextPt.x + x + 0.5f, nextPt.y + y + 0.5f); + work_type J_val = Tex_J::read(nextPt.x + x + 0.5f, nextPt.y + y + 0.5f); work_type diff = J_val - I_val; @@ -309,22 +536,352 @@ namespace pyrlk } } - 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) + // Kernel, uses non texture fetches + template + __global__ void sparseKernel_(Ptr2D I, Ptr2D J, const float2* prevPts, float2* nextPts, uchar* status, float* err, const int level, const int rows, const int cols) { - dim3 grid(ptcount); +#if __CUDA_ARCH__ <= 110 + const int BLOCK_SIZE = 128; +#else + const int BLOCK_SIZE = 256; +#endif - if (level == 0 && err) - sparseKernel<<>>(prevPts, nextPts, status, err, level, rows, cols); - else - sparseKernel<<>>(prevPts, nextPts, status, err, level, rows, cols); + __shared__ float smem1[BLOCK_SIZE]; + __shared__ float smem2[BLOCK_SIZE]; + __shared__ float smem3[BLOCK_SIZE]; - cudaSafeCall( cudaGetLastError() ); + 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_halfWin_x; + prevPt.y -= c_halfWin_y; + + // 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_winSize_y; yBase += blockDim.y, ++i) + { + for (int xBase = threadIdx.x, j = 0; xBase < c_winSize_x; xBase += blockDim.x, ++j) + { + float x = prevPt.x + xBase + 0.5f; + float y = prevPt.y + yBase + 0.5f; + + I_patch[i][j] = ToFloat(I(y, x)); + + // Sharr Deriv + + work_type dIdx = 3.0f * I(y - 1, x + 1) + 10.0f * I(y, x + 1) + 3.0f * I(y + 1, x + 1) - + (3.0f * I(y - 1, x - 1) + 10.0f * I(y, x - 1) + 3.0f * I(y + 1 , x - 1)); + + work_type dIdy = 3.0f * I(y + 1, x - 1) + 10.0f * I(y + 1, x) + 3.0f * I(y+1, x + 1) - + (3.0f * I(y - 1, x - 1) + 10.0f * I(y-1, x) + 3.0f * I(y - 1, x + 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 (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_halfWin_x; + nextPt.y -= c_halfWin_y; + + for (int k = 0; k < c_iters; ++k) + { + if (nextPt.x < -c_halfWin_x || nextPt.x >= cols || nextPt.y < -c_halfWin_y || 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_winSize_y; y += blockDim.y, ++i) + { + for (int x = threadIdx.x, j = 0; x < c_winSize_x; x += blockDim.x, ++j) + { + work_type I_val = I_patch[i][j]; + work_type J_val = ToFloat(J(nextPt.y + y + 0.5f, nextPt.x + x + 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_winSize_y; y += blockDim.y, ++i) + { + for (int x = threadIdx.x, j = 0; x < c_winSize_x; x += blockDim.x, ++j) + { + work_type I_val = I_patch[i][j]; + work_type J_val = ToFloat(J(nextPt.y + y + 0.5f, nextPt.x + x + 0.5f)); + + work_type diff = J_val - I_val; + + accum(errval, abs_(diff)); + } + } + + reduce(smem1, errval, tid, plus()); + } + + if (tid == 0) + { + nextPt.x += c_halfWin_x; + nextPt.y += c_halfWin_y; + + nextPts[blockIdx.x] = nextPt; + + if (calcErr) + err[blockIdx.x] = static_cast(errval) / (3 * c_winSize_x * c_winSize_y); + } + } // __global__ void sparseKernel_ + + + template class sparse_caller + { + public: + static void call(PtrStepSz::vec_type> I, PtrStepSz::vec_type> J, 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); + (void)I; + (void)J; + if (level == 0 && err) + sparseKernel <<>>(prevPts, nextPts, status, err, level, rows, cols); + else + sparseKernel <<>>(prevPts, nextPts, status, err, level, rows, cols); + + cudaSafeCall(cudaGetLastError()); + + if (stream == 0) + cudaSafeCall(cudaDeviceSynchronize()); + } + }; + // Specialization to use non texture path because for some reason the texture path keeps failing accuracy tests + template class sparse_caller<1, PATCH_X, PATCH_Y, unsigned short> + { + public: + typedef typename TypeVec::vec_type work_type; + typedef PtrStepSz Ptr2D; + typedef BrdConstant BrdType; + typedef BorderReader Reader; + typedef LinearFilter Filter; + static void call(Ptr2D I, Ptr2D J, 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_ <<>>( + Filter(Reader(I, BrdType(rows, cols))), + Filter(Reader(J, BrdType(rows, cols))), + prevPts, nextPts, status, err, level, rows, cols); + } + else + { + sparseKernel_ <<>>( + Filter(Reader(I, BrdType(rows, cols))), + Filter(Reader(J, BrdType(rows, cols))), + prevPts, nextPts, status, err, level, rows, cols); + } + cudaSafeCall(cudaGetLastError()); + + if (stream == 0) + cudaSafeCall(cudaDeviceSynchronize()); + } + }; + // Specialization for int because the texture path keeps failing + template class sparse_caller<1, PATCH_X, PATCH_Y, int> + { + public: + typedef typename TypeVec::vec_type work_type; + typedef PtrStepSz Ptr2D; + typedef BrdConstant BrdType; + typedef BorderReader Reader; + typedef LinearFilter Filter; + static void call(Ptr2D I, Ptr2D J, 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_ <<>>( + Filter(Reader(I, BrdType(rows, cols))), + Filter(Reader(J, BrdType(rows, cols))), + prevPts, nextPts, status, err, level, rows, cols); + } + else + { + sparseKernel_ <<>>( + Filter(Reader(I, BrdType(rows, cols))), + Filter(Reader(J, BrdType(rows, cols))), + prevPts, nextPts, status, err, level, rows, cols); + } + cudaSafeCall(cudaGetLastError()); + + if (stream == 0) + cudaSafeCall(cudaDeviceSynchronize()); + } + }; + template class sparse_caller<4, PATCH_X, PATCH_Y, int> + { + public: + typedef typename TypeVec::vec_type work_type; + typedef PtrStepSz Ptr2D; + typedef BrdConstant BrdType; + typedef BorderReader Reader; + typedef LinearFilter Filter; + static void call(Ptr2D I, Ptr2D J, 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_ <<>>( + Filter(Reader(I, BrdType(rows, cols))), + Filter(Reader(J, BrdType(rows, cols))), + prevPts, nextPts, status, err, level, rows, cols); + } + else + { + sparseKernel_ <<>>( + Filter(Reader(I, BrdType(rows, cols))), + Filter(Reader(J, BrdType(rows, cols))), + prevPts, nextPts, status, err, level, rows, cols); + } + cudaSafeCall(cudaGetLastError()); + + if (stream == 0) + cudaSafeCall(cudaDeviceSynchronize()); + } + }; + using namespace cv::cuda::device; + template class sparse_caller<3, PATCH_X, PATCH_Y, T> + { + public: + typedef typename TypeVec::vec_type work_type; + typedef PtrStepSz Ptr2D; + typedef BrdConstant BrdType; + typedef BorderReader Reader; + typedef LinearFilter Filter; + static void call(Ptr2D I, Ptr2D J, 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_ <<>>( + Filter(Reader(I, BrdType(rows, cols))), + Filter(Reader(J, BrdType(rows, cols))), + prevPts, nextPts, status, err, level, rows, cols); + } + else + { + sparseKernel_ <<>>( + Filter(Reader(I, BrdType(rows, cols))), + Filter(Reader(J, BrdType(rows, cols))), + prevPts, nextPts, status, err, level, rows, cols); + } + cudaSafeCall(cudaGetLastError()); + + if (stream == 0) + cudaSafeCall(cudaDeviceSynchronize()); + } + }; - if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); - } template __global__ void denseKernel(PtrStepf u, PtrStepf v, const PtrStepf prevU, const PtrStepf prevV, PtrStepf err, const int rows, const int cols) @@ -484,77 +1041,72 @@ namespace pyrlk cudaSafeCall( cudaMemcpyToSymbolAsync(c_iters, &iters, sizeof(int), 0, cudaMemcpyHostToDevice, stream) ); } - 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) + template struct pyrLK_caller { - 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] = + static void sparse(PtrStepSz::vec_type> I, PtrStepSz::vec_type> J, const float2* prevPts, float2* nextPts, uchar* status, float* err, int ptcount, + int level, dim3 block, dim3 patch, cudaStream_t stream) { - {sparse_caller<1, 1, 1>, sparse_caller<1, 2, 1>, sparse_caller<1, 3, 1>, sparse_caller<1, 4, 1>, sparse_caller<1, 5, 1>}, - {sparse_caller<1, 1, 2>, sparse_caller<1, 2, 2>, sparse_caller<1, 3, 2>, sparse_caller<1, 4, 2>, sparse_caller<1, 5, 2>}, - {sparse_caller<1, 1, 3>, sparse_caller<1, 2, 3>, sparse_caller<1, 3, 3>, sparse_caller<1, 4, 3>, sparse_caller<1, 5, 3>}, - {sparse_caller<1, 1, 4>, sparse_caller<1, 2, 4>, sparse_caller<1, 3, 4>, sparse_caller<1, 4, 4>, sparse_caller<1, 5, 4>}, - {sparse_caller<1, 1, 5>, sparse_caller<1, 2, 5>, sparse_caller<1, 3, 5>, sparse_caller<1, 4, 5>, sparse_caller<1, 5, 5>} - }; + typedef void(*func_t)(PtrStepSz::vec_type> I, PtrStepSz::vec_type> J, + int rows, int cols, const float2* prevPts, float2* nextPts, uchar* status, float* err, int ptcount, + int level, dim3 block, cudaStream_t stream); - bindTexture(&tex_If, I); - bindTexture(&tex_Jf, J); + static const func_t funcs[5][5] = + { + { sparse_caller::call, sparse_caller::call, sparse_caller::call, sparse_caller::call, sparse_caller::call }, + { sparse_caller::call, sparse_caller::call, sparse_caller::call, sparse_caller::call, sparse_caller::call }, + { sparse_caller::call, sparse_caller::call, sparse_caller::call, sparse_caller::call, sparse_caller::call }, + { sparse_caller::call, sparse_caller::call, sparse_caller::call, sparse_caller::call, sparse_caller::call }, + { sparse_caller::call, sparse_caller::call, sparse_caller::call, sparse_caller::call, sparse_caller::call } + }; - funcs[patch.y - 1][patch.x - 1](I.rows, I.cols, prevPts, nextPts, status, err, ptcount, - level, block, stream); - } + Tex_I::bindTexture_(I); + Tex_J::bindTexture_(J); - 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) - { - 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] = - { - {sparse_caller<4, 1, 1>, sparse_caller<4, 2, 1>, sparse_caller<4, 3, 1>, sparse_caller<4, 4, 1>, sparse_caller<4, 5, 1>}, - {sparse_caller<4, 1, 2>, sparse_caller<4, 2, 2>, sparse_caller<4, 3, 2>, sparse_caller<4, 4, 2>, sparse_caller<4, 5, 2>}, - {sparse_caller<4, 1, 3>, sparse_caller<4, 2, 3>, sparse_caller<4, 3, 3>, sparse_caller<4, 4, 3>, sparse_caller<4, 5, 3>}, - {sparse_caller<4, 1, 4>, sparse_caller<4, 2, 4>, sparse_caller<4, 3, 4>, sparse_caller<4, 4, 4>, sparse_caller<4, 5, 4>}, - {sparse_caller<4, 1, 5>, sparse_caller<4, 2, 5>, sparse_caller<4, 3, 5>, sparse_caller<4, 4, 5>, sparse_caller<4, 5, 5>} - }; - - bindTexture(&tex_If4, I); - bindTexture(&tex_Jf4, J); - - funcs[patch.y - 1][patch.x - 1](I.rows, I.cols, prevPts, nextPts, status, err, ptcount, - level, block, stream); - } - - 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); - dim3 grid(divUp(I.cols, block.x), divUp(I.rows, block.y)); - - bindTexture(&tex_Ib, I); - bindTexture(&tex_Jf, J); - - int2 halfWin = make_int2((winSize.x - 1) / 2, (winSize.y - 1) / 2); - const int patchWidth = block.x + 2 * halfWin.x; - const int patchHeight = block.y + 2 * halfWin.y; - size_t smem_size = 3 * patchWidth * patchHeight * sizeof(int); - - if (err.data) - { - denseKernel<<>>(u, v, prevU, prevV, err, I.rows, I.cols); - cudaSafeCall( cudaGetLastError() ); + funcs[patch.y - 1][patch.x - 1](I, J, I.rows, I.cols, prevPts, nextPts, status, err, ptcount, + level, block, stream); } - else + static void dense(PtrStepSzb I, PtrStepSz J, PtrStepSzf u, PtrStepSzf v, PtrStepSzf prevU, PtrStepSzf prevV, PtrStepSzf err, int2 winSize, cudaStream_t stream) { - denseKernel<<>>(u, v, prevU, prevV, PtrStepf(), I.rows, I.cols); - cudaSafeCall( cudaGetLastError() ); - } + dim3 block(16, 16); + dim3 grid(divUp(I.cols, block.x), divUp(I.rows, block.y)); + Tex_I<1, uchar>::bindTexture_(I); + Tex_J<1, T>::bindTexture_(J); - if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); - } + int2 halfWin = make_int2((winSize.x - 1) / 2, (winSize.y - 1) / 2); + const int patchWidth = block.x + 2 * halfWin.x; + const int patchHeight = block.y + 2 * halfWin.y; + size_t smem_size = 3 * patchWidth * patchHeight * sizeof(int); + + if (err.data) + { + denseKernel << > >(u, v, prevU, prevV, err, I.rows, I.cols); + cudaSafeCall(cudaGetLastError()); + } + else + { + denseKernel << > >(u, v, prevU, prevV, PtrStepf(), I.rows, I.cols); + cudaSafeCall(cudaGetLastError()); + } + + if (stream == 0) + cudaSafeCall(cudaDeviceSynchronize()); + } + }; + + template class pyrLK_caller; + template class pyrLK_caller; + template class pyrLK_caller; + template class pyrLK_caller; + + template class pyrLK_caller; + template class pyrLK_caller; + template class pyrLK_caller; + template class pyrLK_caller; + + template class pyrLK_caller; + template class pyrLK_caller; + template class pyrLK_caller; + template class pyrLK_caller; } -#endif /* CUDA_DISABLER */ +#endif /* CUDA_DISABLER */ \ No newline at end of file diff --git a/modules/cudaoptflow/src/precomp.hpp b/modules/cudaoptflow/src/precomp.hpp index 3c818dd4e..d5ac49334 100644 --- a/modules/cudaoptflow/src/precomp.hpp +++ b/modules/cudaoptflow/src/precomp.hpp @@ -52,7 +52,7 @@ #include "opencv2/video.hpp" #include "opencv2/core/private.cuda.hpp" - +#include "opencv2/core/cuda/vec_traits.hpp" #include "opencv2/opencv_modules.hpp" #ifdef HAVE_OPENCV_CUDALEGACY diff --git a/modules/cudaoptflow/src/pyrlk.cpp b/modules/cudaoptflow/src/pyrlk.cpp index 9d7db0a43..dcfd1f66d 100644 --- a/modules/cudaoptflow/src/pyrlk.cpp +++ b/modules/cudaoptflow/src/pyrlk.cpp @@ -56,14 +56,20 @@ Ptr cv::cuda::DensePyrLKOpticalFlow::create(Size, int, in namespace pyrlk { void loadConstants(int2 winSize, int iters, cudaStream_t stream); + template struct pyrLK_caller + { + static void sparse(PtrStepSz::vec_type> I, PtrStepSz::vec_type> J, const float2* prevPts, float2* nextPts, uchar* status, float* err, int ptcount, + int level, dim3 block, dim3 patch, cudaStream_t stream); - 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); - 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); + static void dense(PtrStepSzb I, PtrStepSzf J, PtrStepSzf u, PtrStepSzf v, PtrStepSzf prevU, PtrStepSzf prevV, + PtrStepSzf err, int2 winSize, cudaStream_t stream); + }; - void dense(PtrStepSzb I, PtrStepSzf J, PtrStepSzf u, PtrStepSzf v, PtrStepSzf prevU, PtrStepSzf prevV, - PtrStepSzf err, int2 winSize, cudaStream_t stream); + template void dispatcher(GpuMat I, GpuMat J, const float2* prevPts, float2* nextPts, uchar* status, float* err, int ptcount, + int level, dim3 block, dim3 patch, cudaStream_t stream) + { + pyrLK_caller::sparse(I, J, prevPts, nextPts, status, err, ptcount, level, block, patch, stream); + } } namespace @@ -76,6 +82,9 @@ namespace void sparse(const GpuMat& prevImg, const GpuMat& nextImg, const GpuMat& prevPts, GpuMat& nextPts, GpuMat& status, GpuMat* err, Stream& stream); + void sparse(std::vector& prevPyr, std::vector& nextPyr, const GpuMat& prevPts, GpuMat& nextPts, + GpuMat& status, GpuMat* err, Stream& stream); + void dense(const GpuMat& prevImg, const GpuMat& nextImg, GpuMat& u, GpuMat& v, Stream& stream); protected: @@ -83,8 +92,9 @@ namespace int maxLevel_; int iters_; bool useInitialFlow_; - + void buildImagePyramid(const GpuMat& prevImg, std::vector& prevPyr, const GpuMat& nextImg, std::vector& nextPyr, Stream stream); private: + friend class SparsePyrLKOpticalFlowImpl; std::vector prevPyr_; std::vector nextPyr_; }; @@ -113,6 +123,88 @@ namespace block.z = patch.z = 1; } + void PyrLKOpticalFlowBase::buildImagePyramid(const GpuMat& prevImg, std::vector& prevPyr, const GpuMat& nextImg, std::vector& nextPyr, Stream stream) + { + prevPyr.resize(maxLevel_ + 1); + nextPyr.resize(maxLevel_ + 1); + + int cn = prevImg.channels(); + + CV_Assert(cn == 1 || cn == 3 || cn == 4); + + prevPyr[0] = prevImg; + nextPyr[0] = nextImg; + + for (int level = 1; level <= maxLevel_; ++level) + { + cuda::pyrDown(prevPyr[level - 1], prevPyr[level], stream); + cuda::pyrDown(nextPyr[level - 1], nextPyr[level], stream); + } + } + void PyrLKOpticalFlowBase::sparse(std::vector& prevPyr, std::vector& nextPyr, const GpuMat& prevPts, GpuMat& nextPts, + GpuMat& status, GpuMat* err, Stream& stream) + { + CV_Assert(prevPyr.size() && nextPyr.size() && "Pyramid needs to at least contain the original matrix as the first element"); + CV_Assert(prevPyr[0].size() == nextPyr[0].size()); + CV_Assert(prevPts.rows == 1 && prevPts.type() == CV_32FC2); + CV_Assert(maxLevel_ >= 0); + CV_Assert(winSize_.width > 2 && winSize_.height > 2); + if (useInitialFlow_) + CV_Assert(nextPts.size() == prevPts.size() && nextPts.type() == prevPts.type()); + else + ensureSizeIsEnough(1, prevPts.cols, prevPts.type(), nextPts); + + GpuMat temp1 = (useInitialFlow_ ? nextPts : prevPts).reshape(1); + GpuMat temp2 = nextPts.reshape(1); + cuda::multiply(temp1, Scalar::all(1.0 / (1 << maxLevel_) / 2.0), temp2, 1, -1, stream); + + + ensureSizeIsEnough(1, prevPts.cols, CV_8UC1, status); + status.setTo(Scalar::all(1), stream); + + if (err) + ensureSizeIsEnough(1, prevPts.cols, CV_32FC1, *err); + + if (prevPyr.size() != size_t(maxLevel_ + 1) || nextPyr.size() != size_t(maxLevel_ + 1)) + { + buildImagePyramid(prevPyr[0], prevPyr, nextPyr[0], nextPyr, stream); + } + + dim3 block, patch; + calcPatchSize(winSize_, block, patch); + CV_Assert(patch.x > 0 && patch.x < 6 && patch.y > 0 && patch.y < 6); + pyrlk::loadConstants(make_int2(winSize_.width, winSize_.height), iters_, StreamAccessor::getStream(stream)); + + const int cn = prevPyr[0].channels(); + const int type = prevPyr[0].depth(); + + typedef void(*func_t)(GpuMat I, GpuMat J, const float2* prevPts, float2* nextPts, uchar* status, float* err, int ptcount, + int level, dim3 block, dim3 patch, cudaStream_t stream); + + // Current int datatype is disabled due to pyrDown not implementing it + // while ushort does work, it has significantly worse performance, and thus doesn't pass accuracy tests. + static const func_t funcs[6][4] = + { + { pyrlk::dispatcher , /*pyrlk::dispatcher*/ 0, pyrlk::dispatcher , pyrlk::dispatcher }, + { /*pyrlk::dispatcher*/ 0, /*pyrlk::dispatcher*/ 0, /*pyrlk::dispatcher*/ 0, /*pyrlk::dispatcher*/ 0 }, + { pyrlk::dispatcher , /*pyrlk::dispatcher*/0, pyrlk::dispatcher , pyrlk::dispatcher }, + { /*pyrlk::dispatcher*/ 0, /*pyrlk::dispatcher*/ 0, /*pyrlk::dispatcher*/ 0, /*pyrlk::dispatcher*/0 }, + { pyrlk::dispatcher , /*pyrlk::dispatcher*/ 0, pyrlk::dispatcher , pyrlk::dispatcher }, + { pyrlk::dispatcher , /*pyrlk::dispatcher*/ 0, pyrlk::dispatcher , pyrlk::dispatcher } + }; + + func_t func = funcs[type][cn-1]; + CV_Assert(func != NULL && "Datatype not implemented"); + for (int level = maxLevel_; level >= 0; level--) + { + func(prevPyr[level], nextPyr[level], + prevPts.ptr(), nextPts.ptr(), + status.ptr(), level == 0 && err ? err->ptr() : 0, + prevPts.cols, level, block, patch, + StreamAccessor::getStream(stream)); + } + } + void PyrLKOpticalFlowBase::sparse(const GpuMat& prevImg, const GpuMat& nextImg, const GpuMat& prevPts, GpuMat& nextPts, GpuMat& status, GpuMat* err, Stream& stream) { if (prevPts.empty()) @@ -122,86 +214,14 @@ namespace 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() == prevPts.type() ); - else - ensureSizeIsEnough(1, prevPts.cols, prevPts.type(), nextPts); - - GpuMat temp1 = (useInitialFlow_ ? nextPts : prevPts).reshape(1); - GpuMat temp2 = nextPts.reshape(1); - cuda::multiply(temp1, Scalar::all(1.0 / (1 << maxLevel_) / 2.0), temp2, 1, -1, stream); - - ensureSizeIsEnough(1, prevPts.cols, CV_8UC1, status); - status.setTo(Scalar::all(1), stream); - - if (err) - ensureSizeIsEnough(1, prevPts.cols, CV_32FC1, *err); // build the image pyramids. + buildImagePyramid(prevImg, prevPyr_, nextImg, nextPyr_, stream); - BufferPool pool(stream); + sparse(prevPyr_, nextPyr_, prevPts, nextPts, status, err, stream); - prevPyr_.resize(maxLevel_ + 1); - nextPyr_.resize(maxLevel_ + 1); - - int cn = prevImg.channels(); - - if (cn == 1 || cn == 4) - { - prevImg.convertTo(prevPyr_[0], CV_32F, stream); - nextImg.convertTo(nextPyr_[0], CV_32F, stream); - } - else - { - GpuMat buf = pool.getBuffer(prevImg.size(), CV_MAKE_TYPE(prevImg.depth(), 4)); - - cuda::cvtColor(prevImg, buf, COLOR_BGR2BGRA, 0, stream); - buf.convertTo(prevPyr_[0], CV_32F, stream); - - cuda::cvtColor(nextImg, buf, COLOR_BGR2BGRA, 0, stream); - buf.convertTo(nextPyr_[0], CV_32F, stream); - } - - for (int level = 1; level <= maxLevel_; ++level) - { - cuda::pyrDown(prevPyr_[level - 1], prevPyr_[level], stream); - cuda::pyrDown(nextPyr_[level - 1], nextPyr_[level], stream); - } - - pyrlk::loadConstants(make_int2(winSize_.width, winSize_.height), iters_, StreamAccessor::getStream(stream)); - - for (int level = maxLevel_; level >= 0; level--) - { - if (cn == 1) - { - pyrlk::sparse1(prevPyr_[level], nextPyr_[level], - prevPts.ptr(), nextPts.ptr(), - status.ptr(), - level == 0 && err ? err->ptr() : 0, prevPts.cols, - level, block, patch, - StreamAccessor::getStream(stream)); - } - else - { - pyrlk::sparse4(prevPyr_[level], nextPyr_[level], - prevPts.ptr(), nextPts.ptr(), - status.ptr(), - level == 0 && err ? err->ptr() : 0, prevPts.cols, - level, block, patch, - StreamAccessor::getStream(stream)); - } - } } void PyrLKOpticalFlowBase::dense(const GpuMat& prevImg, const GpuMat& nextImg, GpuMat& u, GpuMat& v, Stream& stream) @@ -250,7 +270,7 @@ namespace { int idx2 = (idx + 1) & 1; - pyrlk::dense(prevPyr_[level], nextPyr_[level], + pyrlk::pyrLK_caller::dense(prevPyr_[level], nextPyr_[level], uPyr[idx], vPyr[idx], uPyr[idx2], vPyr[idx2], PtrStepSzf(), winSize2i, StreamAccessor::getStream(stream)); @@ -289,14 +309,23 @@ namespace OutputArray _err, Stream& stream) { - const GpuMat prevImg = _prevImg.getGpuMat(); - const GpuMat nextImg = _nextImg.getGpuMat(); const GpuMat prevPts = _prevPts.getGpuMat(); GpuMat& nextPts = _nextPts.getGpuMatRef(); GpuMat& status = _status.getGpuMatRef(); GpuMat* err = _err.needed() ? &(_err.getGpuMatRef()) : NULL; - - sparse(prevImg, nextImg, prevPts, nextPts, status, err, stream); + if (_prevImg.kind() == _InputArray::STD_VECTOR_CUDA_GPU_MAT && _prevImg.kind() == _InputArray::STD_VECTOR_CUDA_GPU_MAT) + { + std::vector prevPyr, nextPyr; + _prevImg.getGpuMatVector(prevPyr); + _nextImg.getGpuMatVector(nextPyr); + sparse(prevPyr, nextPyr, prevPts, nextPts, status, err, stream); + } + else + { + const GpuMat prevImg = _prevImg.getGpuMat(); + const GpuMat nextImg = _nextImg.getGpuMat(); + sparse(prevImg, nextImg, prevPts, nextPts, status, err, stream); + } } }; @@ -347,4 +376,4 @@ Ptr cv::cuda::DensePyrLKOpticalFlow::create(Size winSize, return makePtr(winSize, maxLevel, iters, useInitialFlow); } -#endif /* !defined (HAVE_CUDA) */ +#endif /* !defined (HAVE_CUDA) */ \ No newline at end of file diff --git a/modules/cudaoptflow/test/test_optflow.cpp b/modules/cudaoptflow/test/test_optflow.cpp index 63bc461bb..9a3e3e57f 100644 --- a/modules/cudaoptflow/test/test_optflow.cpp +++ b/modules/cudaoptflow/test/test_optflow.cpp @@ -167,33 +167,34 @@ INSTANTIATE_TEST_CASE_P(CUDA_OptFlow, BroxOpticalFlow, ALL_DEVICES); namespace { - IMPLEMENT_PARAM_CLASS(UseGray, bool) + IMPLEMENT_PARAM_CLASS(Chan, int) + IMPLEMENT_PARAM_CLASS(DataType, int) } -PARAM_TEST_CASE(PyrLKOpticalFlow, cv::cuda::DeviceInfo, UseGray) +PARAM_TEST_CASE(PyrLKOpticalFlow, cv::cuda::DeviceInfo, Chan, DataType) { cv::cuda::DeviceInfo devInfo; - bool useGray; - + int channels; + int dataType; virtual void SetUp() { devInfo = GET_PARAM(0); - useGray = GET_PARAM(1); - + channels = GET_PARAM(1); + dataType = GET_PARAM(2); cv::cuda::setDevice(devInfo.deviceID()); } }; CUDA_TEST_P(PyrLKOpticalFlow, Sparse) { - cv::Mat frame0 = readImage("opticalflow/frame0.png", useGray ? cv::IMREAD_GRAYSCALE : cv::IMREAD_COLOR); + cv::Mat frame0 = readImage("opticalflow/frame0.png", channels == 1 ? cv::IMREAD_GRAYSCALE : cv::IMREAD_COLOR); ASSERT_FALSE(frame0.empty()); - cv::Mat frame1 = readImage("opticalflow/frame1.png", useGray ? cv::IMREAD_GRAYSCALE : cv::IMREAD_COLOR); + cv::Mat frame1 = readImage("opticalflow/frame1.png", channels == 1 ? cv::IMREAD_GRAYSCALE : cv::IMREAD_COLOR); ASSERT_FALSE(frame1.empty()); cv::Mat gray_frame; - if (useGray) + if (channels == 1) gray_frame = frame0; else cv::cvtColor(frame0, gray_frame, cv::COLOR_BGR2GRAY); @@ -208,22 +209,32 @@ CUDA_TEST_P(PyrLKOpticalFlow, Sparse) cv::Ptr pyrLK = cv::cuda::SparsePyrLKOpticalFlow::create(); - cv::cuda::GpuMat d_nextPts; - cv::cuda::GpuMat d_status; - pyrLK->calc(loadMat(frame0), loadMat(frame1), d_pts, d_nextPts, d_status); - - std::vector nextPts(d_nextPts.cols); - cv::Mat nextPts_mat(1, d_nextPts.cols, CV_32FC2, (void*) &nextPts[0]); - d_nextPts.download(nextPts_mat); - - std::vector status(d_status.cols); - cv::Mat status_mat(1, d_status.cols, CV_8UC1, (void*) &status[0]); - d_status.download(status_mat); - std::vector nextPts_gold; std::vector status_gold; cv::calcOpticalFlowPyrLK(frame0, frame1, pts, nextPts_gold, status_gold, cv::noArray()); + + cv::cuda::GpuMat d_nextPts; + cv::cuda::GpuMat d_status; + cv::Mat converted0, converted1; + if(channels == 4) + { + cv::cvtColor(frame0, frame0, cv::COLOR_BGR2BGRA); + cv::cvtColor(frame1, frame1, cv::COLOR_BGR2BGRA); + } + frame0.convertTo(converted0, dataType); + frame1.convertTo(converted1, dataType); + + pyrLK->calc(loadMat(converted0), loadMat(converted1), d_pts, d_nextPts, d_status); + + std::vector nextPts(d_nextPts.cols); + cv::Mat nextPts_mat(1, d_nextPts.cols, CV_32FC2, (void*)&nextPts[0]); + d_nextPts.download(nextPts_mat); + + std::vector status(d_status.cols); + cv::Mat status_mat(1, d_status.cols, CV_8UC1, (void*)&status[0]); + d_status.download(status_mat); + ASSERT_EQ(nextPts_gold.size(), nextPts.size()); ASSERT_EQ(status_gold.size(), status.size()); @@ -251,11 +262,16 @@ CUDA_TEST_P(PyrLKOpticalFlow, Sparse) double bad_ratio = static_cast(mistmatch) / nextPts.size(); ASSERT_LE(bad_ratio, 0.01); + + } INSTANTIATE_TEST_CASE_P(CUDA_OptFlow, PyrLKOpticalFlow, testing::Combine( ALL_DEVICES, - testing::Values(UseGray(true), UseGray(false)))); + testing::Values(Chan(1), Chan(3), Chan(4)), + testing::Values(DataType(CV_8U), DataType(CV_16U), DataType(CV_32S), DataType(CV_32F)))); + + ////////////////////////////////////////////////////// // FarnebackOpticalFlow @@ -385,4 +401,4 @@ INSTANTIATE_TEST_CASE_P(CUDA_OptFlow, OpticalFlowDual_TVL1, testing::Combine( ALL_DEVICES, testing::Values(Gamma(0.0), Gamma(1.0)))); -#endif // HAVE_CUDA +#endif // HAVE_CUDA \ No newline at end of file diff --git a/modules/cudawarping/src/cuda/pyr_down.cu b/modules/cudawarping/src/cuda/pyr_down.cu index 3207d65cb..03e791dcf 100644 --- a/modules/cudawarping/src/cuda/pyr_down.cu +++ b/modules/cudawarping/src/cuda/pyr_down.cu @@ -212,10 +212,10 @@ namespace cv { namespace cuda { namespace device template void pyrDown_gpu(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); template void pyrDown_gpu(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); - //template void pyrDown_gpu(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); + template void pyrDown_gpu(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); //template void pyrDown_gpu(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); - //template void pyrDown_gpu(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); - //template void pyrDown_gpu(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); + template void pyrDown_gpu(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); + template void pyrDown_gpu(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); template void pyrDown_gpu(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); //template void pyrDown_gpu(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); @@ -225,4 +225,4 @@ namespace cv { namespace cuda { namespace device }}} // namespace cv { namespace cuda { namespace cudev -#endif /* CUDA_DISABLER */ +#endif /* CUDA_DISABLER */ \ No newline at end of file diff --git a/modules/cudawarping/src/pyramids.cpp b/modules/cudawarping/src/pyramids.cpp index 0cb0f5de5..817a16715 100644 --- a/modules/cudawarping/src/pyramids.cpp +++ b/modules/cudawarping/src/pyramids.cpp @@ -74,7 +74,7 @@ void cv::cuda::pyrDown(InputArray _src, OutputArray _dst, Stream& stream) {0 /*pyrDown_gpu*/, 0 /*pyrDown_gpu*/ , 0 /*pyrDown_gpu*/, 0 /*pyrDown_gpu*/}, {pyrDown_gpu , 0 /*pyrDown_gpu*/, pyrDown_gpu , pyrDown_gpu }, {pyrDown_gpu , 0 /*pyrDown_gpu*/ , pyrDown_gpu , pyrDown_gpu }, - {0 /*pyrDown_gpu*/ , 0 /*pyrDown_gpu*/ , 0 /*pyrDown_gpu*/ , 0 /*pyrDown_gpu*/ }, + {pyrDown_gpu , 0 /*pyrDown_gpu*/ , pyrDown_gpu , pyrDown_gpu }, {pyrDown_gpu , 0 /*pyrDown_gpu*/ , pyrDown_gpu , pyrDown_gpu } }; @@ -131,4 +131,4 @@ void cv::cuda::pyrUp(InputArray _src, OutputArray _dst, Stream& stream) func(src, dst, StreamAccessor::getStream(stream)); } -#endif +#endif \ No newline at end of file