mirror/opencv
1
0
Fork 0
mirror of https://github.com/opencv/opencv.git synced 2025-06-07 09:25:45 +08:00
Code Issues Actions Packages Projects Releases Wiki Activity

reduce filter2d instantiates for tiny build

Browse Source
This commit is contained in:
Vladislav Vinogradov 2015-03-04 14:24:18 +03:00
parent 00c36e88ef
commit 8e49ab1d3b
2 changed files with 20 additions and 1 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

19
modules/gpu/src/cuda/imgproc.cu
View File

@ -985,6 +985,16 @@ namespace cv { namespace gpu { namespace device
int borderMode, const float* borderValue, cudaStream_t stream)
{
typedef void (*func_t)(const PtrStepSz<T> srcWhole, int xoff, int yoff, PtrStepSz<D> dst, int kWidth, int kHeight, int anchorX, int anchorY, const float* borderValue, cudaStream_t stream);
#ifdef OPENCV_TINY_GPU_MODULE
static const func_t funcs[] =
{
Filter2DCaller<T, D, BrdReflect101>::call,
Filter2DCaller<T, D, BrdReplicate>::call,
Filter2DCaller<T, D, BrdConstant>::call,
Filter2DCaller<T, D, BrdReflect>::call,
0
};
#else
static const func_t funcs[] =
{
Filter2DCaller<T, D, BrdReflect101>::call,
@ -993,19 +1003,26 @@ namespace cv { namespace gpu { namespace device
Filter2DCaller<T, D, BrdReflect>::call,
Filter2DCaller<T, D, BrdWrap>::call
};
#endif
const func_t func = funcs[borderMode];
if (!func)
cv::gpu::error("Unsupported input parameters for filter2D", __FILE__, __LINE__, "");
if (stream == 0)
cudaSafeCall( cudaMemcpyToSymbol(c_filter2DKernel, kernel, kWidth * kHeight * sizeof(float), 0, cudaMemcpyDeviceToDevice) );
else
cudaSafeCall( cudaMemcpyToSymbolAsync(c_filter2DKernel, kernel, kWidth * kHeight * sizeof(float), 0, cudaMemcpyDeviceToDevice, stream) );
funcs[borderMode](static_cast< PtrStepSz<T> >(srcWhole), ofsX, ofsY, static_cast< PtrStepSz<D> >(dst), kWidth, kHeight, anchorX, anchorY, borderValue, stream);
func(static_cast< PtrStepSz<T> >(srcWhole), ofsX, ofsY, static_cast< PtrStepSz<D> >(dst), kWidth, kHeight, anchorX, anchorY, borderValue, stream);
}
template void filter2D_gpu<uchar, uchar>(PtrStepSzb srcWhole, int ofsX, int ofsY, PtrStepSzb dst, int kWidth, int kHeight, int anchorX, int anchorY, const float* kernel, int borderMode, const float* borderValue, cudaStream_t stream);
template void filter2D_gpu<uchar4, uchar4>(PtrStepSzb srcWhole, int ofsX, int ofsY, PtrStepSzb dst, int kWidth, int kHeight, int anchorX, int anchorY, const float* kernel, int borderMode, const float* borderValue, cudaStream_t stream);
#ifndef OPENCV_TINY_GPU_MODULE
template void filter2D_gpu<ushort, ushort>(PtrStepSzb srcWhole, int ofsX, int ofsY, PtrStepSzb dst, int kWidth, int kHeight, int anchorX, int anchorY, const float* kernel, int borderMode, const float* borderValue, cudaStream_t stream);
template void filter2D_gpu<ushort4, ushort4>(PtrStepSzb srcWhole, int ofsX, int ofsY, PtrStepSzb dst, int kWidth, int kHeight, int anchorX, int anchorY, const float* kernel, int borderMode, const float* borderValue, cudaStream_t stream);
#endif
template void filter2D_gpu<float, float>(PtrStepSzb srcWhole, int ofsX, int ofsY, PtrStepSzb dst, int kWidth, int kHeight, int anchorX, int anchorY, const float* kernel, int borderMode, const float* borderValue, cudaStream_t stream);
template void filter2D_gpu<float4, float4>(PtrStepSzb srcWhole, int ofsX, int ofsY, PtrStepSzb dst, int kWidth, int kHeight, int anchorX, int anchorY, const float* kernel, int borderMode, const float* borderValue, cudaStream_t stream);
} // namespace imgproc

2
modules/gpu/src/filtering.cpp
View File

@ -789,12 +789,14 @@ Ptr<BaseFilter_GPU> cv::gpu::getLinearFilter_GPU(int srcType, int dstType, const
case CV_8UC4:
func = filter2D_gpu<uchar4, uchar4>;
break;
#ifndef OPENCV_TINY_GPU_MODULE
case CV_16UC1:
func = filter2D_gpu<ushort, ushort>;
break;
case CV_16UC4:
func = filter2D_gpu<ushort4, ushort4>;
break;
#endif
case CV_32FC1:
func = filter2D_gpu<float, float>;
break;