CUDA SeparableLinearFilter might fail because of too large grid dimensions
Author: PierreChatelierCreated Jun 19, 2026Updated Sep 10, 2026
OpenCV 4.14.0 (beta)
A CUDA grid size is usually limited by ($2^{31}-1$, 65535, 65535) https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/compute-capabilities.html#features-and-technical-specifications It is not taken into account in OpenCV CUDA SeparableLinearFilter (at least) and some calls might become wrong.
The problem lies in row_filter.hpp
const dim3 block(BLOCK_DIM_X, BLOCK_DIM_Y);
const dim3 grid(divUp(src.cols, BLOCK_DIM_X * PATCH_PER_BLOCK), divUp(src.rows, BLOCK_DIM_Y));//SHOULD BE BOUNDED !!!
B<T> brd(src.cols);
linearRowFilter<KSIZE, T, D><<<grid, block, 0, stream>>>(src, dst, kernel, anchor, brd);
cudaSafeCall( cudaGetLastError() );Easy reproducer :
const cv::Size spatialSize(3, 65535*8);//OK, will generate a grid size of (1, 65535, 1)
cv::cuda::GpuMat src = cv::cuda::createContinuous(spatialSize, CV_32FC1);
cv::cuda::GpuMat dst = cv::cuda::createContinuous(spatialSize, CV_32FC1);
cv::Mat kernel = (cv::Mat_<float>(1, 3) << -1, 0, 1);
cv::Ptr<cv::cuda::Filter> filter = cv::cuda::createSeparableLinearFilter(src.type(), dst.type(), kernel, cv::noArray(), cv::Point(-1, -1), cv::BORDER_REPLICATE, cv::BORDER_REPLICATE);
filter->apply(src, dst);const cv::Size spatialSize(3, 65536*8);//WRONG, will generate a grid size of (1, 65536, 1)
cv::cuda::GpuMat src = cv::cuda::createContinuous(spatialSize, CV_32FC1);
cv::cuda::GpuMat dst = cv::cuda::createContinuous(spatialSize, CV_32FC1);
cv::Mat kernel = (cv::Mat_<float>(1, 3) << -1, 0, 1);
cv::Ptr<cv::cuda::Filter> filter = cv::cuda::createSeparableLinearFilter(src.type(), dst.type(), kernel, cv::noArray(), cv::Point(-1, -1), cv::BORDER_REPLICATE, cv::BORDER_REPLICATE);
filter->apply(src, dst); //<-exception because of CUDA failureSource: opencv/opencv_contrib