From 8c08714b8c5dc622503076d9c454f86e3c7d1da4 Mon Sep 17 00:00:00 2001 From: Alexander Karsakov Date: Thu, 11 Sep 2014 18:11:23 +0400 Subject: [PATCH] Remove two "set" kernel call --- modules/imgproc/src/featureselect.cpp | 9 +++++---- modules/imgproc/src/opencl/gftt.cl | 4 +++- 2 files changed, 8 insertions(+), 5 deletions(-) diff --git a/modules/imgproc/src/featureselect.cpp b/modules/imgproc/src/featureselect.cpp index 769624cb9b..c61ceb013a 100644 --- a/modules/imgproc/src/featureselect.cpp +++ b/modules/imgproc/src/featureselect.cpp @@ -85,6 +85,7 @@ static bool ocl_goodFeaturesToTrack( InputArray _image, OutputArray _corners, size_t total, i, j, ncorners = 0, possibleCornersCount = std::max(1024, static_cast(imgsize.area() * 0.1)); bool haveMask = !_mask.empty(); + UMat counter(1, 1, CV_32SC1); // find threshold { @@ -108,7 +109,8 @@ static bool ocl_goodFeaturesToTrack( InputArray _image, OutputArray _corners, ocl::KernelArg eigarg = ocl::KernelArg::ReadOnlyNoSize(eig), dbarg = ocl::KernelArg::PtrWriteOnly(maxEigenValue), - maskarg = ocl::KernelArg::ReadOnlyNoSize(mask); + maskarg = ocl::KernelArg::ReadOnlyNoSize(mask), + counterarg = ocl::KernelArg::PtrReadWrite(counter); if (haveMask) k.args(eigarg, eig.cols, (int)eig.total(), dbarg, maskarg); @@ -125,7 +127,7 @@ static bool ocl_goodFeaturesToTrack( InputArray _image, OutputArray _corners, if (k2.empty()) return false; - k2.args(dbarg, (float)qualityLevel); + k2.args(dbarg, (float)qualityLevel, counterarg); if (!k2.runTask(false)) return false; @@ -138,8 +140,7 @@ static bool ocl_goodFeaturesToTrack( InputArray _image, OutputArray _corners, if (k.empty()) return false; - UMat counter(1, 1, CV_32SC1, Scalar::all(0)), - corners(1, (int)possibleCornersCount, CV_32FC2, Scalar::all(-1)); + UMat corners(1, (int)possibleCornersCount, CV_32FC2); CV_Assert(sizeof(Corner) == corners.elemSize()); ocl::KernelArg eigarg = ocl::KernelArg::ReadOnlyNoSize(eig), diff --git a/modules/imgproc/src/opencl/gftt.cl b/modules/imgproc/src/opencl/gftt.cl index a63ad044b1..0ecbf3a31b 100644 --- a/modules/imgproc/src/opencl/gftt.cl +++ b/modules/imgproc/src/opencl/gftt.cl @@ -91,7 +91,8 @@ __kernel void maxEigenVal(__global const uchar * srcptr, int src_step, int src_o *(__global float *)(dstptr + (int)sizeof(float) * gid) = localmem_max[0]; } -__kernel void maxEigenValTask(__global float * dst, float qualityLevel) +__kernel void maxEigenValTask(__global float * dst, float qualityLevel, + __global int * counter) { float maxval = -FLT_MAX; @@ -100,6 +101,7 @@ __kernel void maxEigenValTask(__global float * dst, float qualityLevel) maxval = max(maxval, dst[x]); dst[0] = maxval * qualityLevel; + counter[0] = 0; } #elif OP_FIND_CORNERS