Skip to content

Fix CUDA mem leak in GFTT and move CUDA malloc out of critical path #3324

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 2 commits into from
Aug 12, 2022
Merged

Fix CUDA mem leak in GFTT and move CUDA malloc out of critical path #3324

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
843df49
fix cuda mem leak and move cuda malloc out of critical path
kevinchristensen1 Aug 10, 2022
1d1dbe3
destroy texture object
kevinchristensen1 Aug 10, 2022
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
13 changes: 5 additions & 8 deletions modules/cudaimgproc/src/cuda/gftt.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -87,25 +87,22 @@ namespace cv { namespace cuda { namespace device
}
}

int findCorners_gpu(const cudaTextureObject_t &eigTex, const int &rows, const int &cols, float threshold, PtrStepSzb mask, float2* corners, int max_count, cudaStream_t stream)
int findCorners_gpu(const cudaTextureObject_t &eigTex, const int &rows, const int &cols, float threshold, PtrStepSzb mask, float2* corners, int max_count, int* counterPtr, cudaStream_t stream)
{
int* counter_ptr;
cudaSafeCall( cudaMalloc(&counter_ptr, sizeof(int)) );

cudaSafeCall( cudaMemsetAsync(counter_ptr, 0, sizeof(int), stream) );
cudaSafeCall( cudaMemsetAsync(counterPtr, 0, sizeof(int), stream) );

dim3 block(16, 16);
dim3 grid(divUp(cols, block.x), divUp(rows, block.y));

if (mask.data)
findCorners<<<grid, block, 0, stream>>>(threshold, SingleMask(mask), corners, max_count, rows, cols, eigTex, counter_ptr);
findCorners<<<grid, block, 0, stream>>>(threshold, SingleMask(mask), corners, max_count, rows, cols, eigTex, counterPtr);
else
findCorners<<<grid, block, 0, stream>>>(threshold, WithOutMask(), corners, max_count, rows, cols, eigTex, counter_ptr);
findCorners<<<grid, block, 0, stream>>>(threshold, WithOutMask(), corners, max_count, rows, cols, eigTex, counterPtr);

cudaSafeCall( cudaGetLastError() );

int count;
cudaSafeCall( cudaMemcpyAsync(&count, counter_ptr, sizeof(int), cudaMemcpyDeviceToHost, stream) );
cudaSafeCall( cudaMemcpyAsync(&count, counterPtr, sizeof(int), cudaMemcpyDeviceToHost, stream) );
if (stream)
cudaSafeCall(cudaStreamSynchronize(stream));
else
Expand Down
18 changes: 14 additions & 4 deletions modules/cudaimgproc/src/gftt.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -55,7 +55,7 @@ namespace cv { namespace cuda { namespace device
{
namespace gfft
{
int findCorners_gpu(const cudaTextureObject_t &eigTex_, const int &rows, const int &cols, float threshold, PtrStepSzb mask, float2* corners, int max_count, cudaStream_t stream);
int findCorners_gpu(const cudaTextureObject_t &eigTex_, const int &rows, const int &cols, float threshold, PtrStepSzb mask, float2* corners, int max_count, int* counterPtr, cudaStream_t stream);
void sortCorners_gpu(const cudaTextureObject_t &eigTex_, float2* corners, int count, cudaStream_t stream);
}
}}}
Expand All @@ -67,7 +67,7 @@ namespace
public:
GoodFeaturesToTrackDetector(int srcType, int maxCorners, double qualityLevel, double minDistance,
int blockSize, bool useHarrisDetector, double harrisK);

~GoodFeaturesToTrackDetector();
void detect(InputArray image, OutputArray corners, InputArray mask, Stream& stream);

private:
Expand All @@ -82,6 +82,8 @@ namespace
GpuMat buf_;
GpuMat eig_;
GpuMat tmpCorners_;

int* counterPtr_;
};

GoodFeaturesToTrackDetector::GoodFeaturesToTrackDetector(int srcType, int maxCorners, double qualityLevel, double minDistance,
Expand All @@ -93,6 +95,12 @@ namespace
cornerCriteria_ = useHarrisDetector ?
cuda::createHarrisCorner(srcType, blockSize, 3, harrisK) :
cuda::createMinEigenValCorner(srcType, blockSize, 3);
cudaSafeCall(cudaMalloc(&counterPtr_, sizeof(int)));
}

GoodFeaturesToTrackDetector::~GoodFeaturesToTrackDetector()
{
cudaSafeCall(cudaFree(counterPtr_));
}

void GoodFeaturesToTrackDetector::detect(InputArray _image, OutputArray _corners, InputArray _mask, Stream& stream)
Expand Down Expand Up @@ -125,17 +133,19 @@ namespace
PtrStepSzf eig = eig_;
cv::cuda::device::createTextureObjectPitch2D<float>(&eigTex_, eig, texDesc);

int total = findCorners_gpu(eigTex_, eig_.rows, eig_.cols, static_cast<float>(maxVal * qualityLevel_), mask, tmpCorners_.ptr<float2>(), tmpCorners_.cols, stream_);

int total = findCorners_gpu(eigTex_, eig_.rows, eig_.cols, static_cast<float>(maxVal * qualityLevel_), mask, tmpCorners_.ptr<float2>(), tmpCorners_.cols, counterPtr_, stream_);

if (total == 0)
{
_corners.release();
cudaSafeCall( cudaDestroyTextureObject(eigTex_) );
return;
}

sortCorners_gpu(eigTex_, tmpCorners_.ptr<float2>(), total, stream_);

cudaSafeCall( cudaDestroyTextureObject(eigTex_) );

if (minDistance_ < 1)
{
tmpCorners_.colRange(0, maxCorners_ > 0 ? std::min(maxCorners_, total) : total).copyTo(_corners, stream);
Expand Down