diff --git a/modules/gpu/CMakeLists.txt b/modules/gpu/CMakeLists.txt index 70428a40b9..c5576e6bff 100644 --- a/modules/gpu/CMakeLists.txt +++ b/modules/gpu/CMakeLists.txt @@ -3,7 +3,7 @@ if(ANDROID OR IOS) endif() set(the_description "GPU-accelerated Computer Vision") -ocv_add_module(gpu opencv_imgproc opencv_calib3d opencv_objdetect) +ocv_add_module(gpu opencv_imgproc opencv_calib3d opencv_objdetect opencv_video) ocv_module_include_directories("${CMAKE_CURRENT_SOURCE_DIR}/src/cuda") diff --git a/modules/gpu/perf/perf_video.cpp b/modules/gpu/perf/perf_video.cpp index b349bc54ee..4ba17809ea 100644 --- a/modules/gpu/perf/perf_video.cpp +++ b/modules/gpu/perf/perf_video.cpp @@ -193,4 +193,36 @@ INSTANTIATE_TEST_CASE_P(Video, PyrLKOpticalFlowSparse, testing::Combine testing::Values(17, 21) )); +////////////////////////////////////////////////////// +// PyrLKOpticalFlowDense + +GPU_PERF_TEST_1(PyrLKOpticalFlowDense, cv::gpu::DeviceInfo) +{ + cv::gpu::DeviceInfo devInfo = GetParam(); + + cv::gpu::setDevice(devInfo.deviceID()); + + cv::Mat frame0_host = readImage("gpu/perf/aloe.jpg", cv::IMREAD_GRAYSCALE); + cv::Mat frame1_host = readImage("gpu/perf/aloeR.jpg", cv::IMREAD_GRAYSCALE); + + ASSERT_FALSE(frame0_host.empty()); + ASSERT_FALSE(frame1_host.empty()); + + cv::gpu::GpuMat frame0(frame0_host); + cv::gpu::GpuMat frame1(frame1_host); + cv::gpu::GpuMat u; + cv::gpu::GpuMat v; + + cv::gpu::PyrLKOpticalFlow pyrLK; + + declare.time(10); + + TEST_CYCLE() + { + pyrLK.dense(frame0, frame1, u, v); + } +} + +INSTANTIATE_TEST_CASE_P(Video, PyrLKOpticalFlowDense, ALL_DEVICES); + #endif diff --git a/modules/gpu/src/cuda/surf.cu b/modules/gpu/src/cuda/surf.cu index 67d474dd9e..9d265cc169 100644 --- a/modules/gpu/src/cuda/surf.cu +++ b/modules/gpu/src/cuda/surf.cu @@ -541,8 +541,8 @@ namespace cv { namespace gpu { namespace device angle = atan2f(Y, X); if (angle < 0) - angle += 2.0f * CV_PI; - angle *= 180.0f / CV_PI; + angle += 2.0f * CV_PI_F; + angle *= 180.0f / CV_PI_F; } } s_X[tid] = X; @@ -636,8 +636,8 @@ namespace cv { namespace gpu { namespace device { float kp_dir = atan2f(besty, bestx); if (kp_dir < 0) - kp_dir += 2.0f * CV_PI; - kp_dir *= 180.0f / CV_PI; + kp_dir += 2.0f * CV_PI_F; + kp_dir *= 180.0f / CV_PI_F; featureDir[blockIdx.x] = kp_dir; } @@ -724,7 +724,7 @@ namespace cv { namespace gpu { namespace device const float centerX = featureX[blockIdx.x]; const float centerY = featureY[blockIdx.x]; const float size = featureSize[blockIdx.x]; - const float descriptor_dir = featureDir[blockIdx.x] * (float)(CV_PI / 180); + const float descriptor_dir = featureDir[blockIdx.x] * (float)(CV_PI_F / 180.0f); /* The sampling intervals and wavelet sized for selecting an orientation and building the keypoint descriptor are defined relative to 's' */ @@ -817,6 +817,7 @@ namespace cv { namespace gpu { namespace device calc_dx_dy(sdx, sdy, featureX, featureY, featureSize, featureDir); __syncthreads(); + const int tid = threadIdx.y * blockDim.x + threadIdx.x; if (tid < 25)