namespace cv { namespace gpu
{
- enum
- {
- BORDER_REFLECT101_GPU = 0,
- BORDER_REPLICATE_GPU,
- BORDER_CONSTANT_GPU,
- BORDER_REFLECT_GPU,
- BORDER_WRAP_GPU
- };
-
namespace cudev
{
__host__ __device__ __forceinline__ int divUp(int total, int grain)
namespace cv { namespace gpu {
CV_EXPORTS cv::String getNppErrorMessage(int code);
CV_EXPORTS cv::String getCudaDriverApiErrorMessage(int code);
-
- // Converts CPU border extrapolation mode into GPU internal analogue.
- // Returns true if the GPU analogue exists, false otherwise.
- CV_EXPORTS bool tryConvertToGpuBorderType(int cpuBorderType, int& gpuBorderType);
}}
#ifndef HAVE_CUDA
return getErrorString(code, cu_errors, cu_errors_num);
#endif
}
-
-bool cv::gpu::tryConvertToGpuBorderType(int cpuBorderType, int& gpuBorderType)
-{
-#ifndef HAVE_CUDA
- (void) cpuBorderType;
- (void) gpuBorderType;
- return false;
-#else
- switch (cpuBorderType)
- {
- case IPL_BORDER_REFLECT_101:
- gpuBorderType = cv::gpu::BORDER_REFLECT101_GPU;
- return true;
- case IPL_BORDER_REPLICATE:
- gpuBorderType = cv::gpu::BORDER_REPLICATE_GPU;
- return true;
- case IPL_BORDER_CONSTANT:
- gpuBorderType = cv::gpu::BORDER_CONSTANT_GPU;
- return true;
- case IPL_BORDER_REFLECT:
- gpuBorderType = cv::gpu::BORDER_REFLECT_GPU;
- return true;
- case IPL_BORDER_WRAP:
- gpuBorderType = cv::gpu::BORDER_WRAP_GPU;
- return true;
- default:
- return false;
- };
-#endif
-}
void cv::gpu::copyMakeBorder(const GpuMat& src, GpuMat& dst, int top, int bottom, int left, int right, int borderType, const Scalar& value, Stream& s)
{
CV_Assert(src.depth() <= CV_32F && src.channels() <= 4);
- CV_Assert(borderType == IPL_BORDER_REFLECT_101 || borderType == IPL_BORDER_REPLICATE || borderType == IPL_BORDER_CONSTANT || borderType == IPL_BORDER_REFLECT || borderType == IPL_BORDER_WRAP);
+ CV_Assert(borderType == BORDER_REFLECT_101 || borderType == BORDER_REPLICATE || borderType == BORDER_CONSTANT || borderType == BORDER_REFLECT || borderType == BORDER_WRAP);
dst.create(src.rows + top + bottom, src.cols + left + right, src.type());
cudaStream_t stream = StreamAccessor::getStream(s);
- if (borderType == IPL_BORDER_CONSTANT && (src.type() == CV_8UC1 || src.type() == CV_8UC4 || src.type() == CV_32SC1 || src.type() == CV_32FC1))
+ if (borderType == BORDER_CONSTANT && (src.type() == CV_8UC1 || src.type() == CV_8UC4 || src.type() == CV_32SC1 || src.type() == CV_32FC1))
{
NppiSize srcsz;
srcsz.width = src.cols;
caller_t func = callers[src.depth()][src.channels() - 1];
CV_Assert(func != 0);
- int gpuBorderType;
- CV_Assert(tryConvertToGpuBorderType(borderType, gpuBorderType));
-
- func(src, dst, top, left, gpuBorderType, value, stream);
+ func(src, dst, top, left, borderType, value, stream);
}
}
static const caller_t callers[5] =
{
- CopyMakeBorderDispatcher<BrdReflect101, vec_type>::call,
- CopyMakeBorderDispatcher<BrdReplicate, vec_type>::call,
CopyMakeBorderDispatcher<BrdConstant, vec_type>::call,
+ CopyMakeBorderDispatcher<BrdReplicate, vec_type>::call,
CopyMakeBorderDispatcher<BrdReflect, vec_type>::call,
- CopyMakeBorderDispatcher<BrdWrap, vec_type>::call
+ CopyMakeBorderDispatcher<BrdWrap, vec_type>::call,
+ CopyMakeBorderDispatcher<BrdReflect101, vec_type>::call
};
callers[borderMode](PtrStepSz<vec_type>(src), PtrStepSz<vec_type>(dst), top, left, borderValue, stream);
{
{
0,
- column_filter::caller< 1, T, D, BrdColReflect101>,
- column_filter::caller< 2, T, D, BrdColReflect101>,
- column_filter::caller< 3, T, D, BrdColReflect101>,
- column_filter::caller< 4, T, D, BrdColReflect101>,
- column_filter::caller< 5, T, D, BrdColReflect101>,
- column_filter::caller< 6, T, D, BrdColReflect101>,
- column_filter::caller< 7, T, D, BrdColReflect101>,
- column_filter::caller< 8, T, D, BrdColReflect101>,
- column_filter::caller< 9, T, D, BrdColReflect101>,
- column_filter::caller<10, T, D, BrdColReflect101>,
- column_filter::caller<11, T, D, BrdColReflect101>,
- column_filter::caller<12, T, D, BrdColReflect101>,
- column_filter::caller<13, T, D, BrdColReflect101>,
- column_filter::caller<14, T, D, BrdColReflect101>,
- column_filter::caller<15, T, D, BrdColReflect101>,
- column_filter::caller<16, T, D, BrdColReflect101>,
- column_filter::caller<17, T, D, BrdColReflect101>,
- column_filter::caller<18, T, D, BrdColReflect101>,
- column_filter::caller<19, T, D, BrdColReflect101>,
- column_filter::caller<20, T, D, BrdColReflect101>,
- column_filter::caller<21, T, D, BrdColReflect101>,
- column_filter::caller<22, T, D, BrdColReflect101>,
- column_filter::caller<23, T, D, BrdColReflect101>,
- column_filter::caller<24, T, D, BrdColReflect101>,
- column_filter::caller<25, T, D, BrdColReflect101>,
- column_filter::caller<26, T, D, BrdColReflect101>,
- column_filter::caller<27, T, D, BrdColReflect101>,
- column_filter::caller<28, T, D, BrdColReflect101>,
- column_filter::caller<29, T, D, BrdColReflect101>,
- column_filter::caller<30, T, D, BrdColReflect101>,
- column_filter::caller<31, T, D, BrdColReflect101>,
- column_filter::caller<32, T, D, BrdColReflect101>
+ column_filter::caller< 1, T, D, BrdColConstant>,
+ column_filter::caller< 2, T, D, BrdColConstant>,
+ column_filter::caller< 3, T, D, BrdColConstant>,
+ column_filter::caller< 4, T, D, BrdColConstant>,
+ column_filter::caller< 5, T, D, BrdColConstant>,
+ column_filter::caller< 6, T, D, BrdColConstant>,
+ column_filter::caller< 7, T, D, BrdColConstant>,
+ column_filter::caller< 8, T, D, BrdColConstant>,
+ column_filter::caller< 9, T, D, BrdColConstant>,
+ column_filter::caller<10, T, D, BrdColConstant>,
+ column_filter::caller<11, T, D, BrdColConstant>,
+ column_filter::caller<12, T, D, BrdColConstant>,
+ column_filter::caller<13, T, D, BrdColConstant>,
+ column_filter::caller<14, T, D, BrdColConstant>,
+ column_filter::caller<15, T, D, BrdColConstant>,
+ column_filter::caller<16, T, D, BrdColConstant>,
+ column_filter::caller<17, T, D, BrdColConstant>,
+ column_filter::caller<18, T, D, BrdColConstant>,
+ column_filter::caller<19, T, D, BrdColConstant>,
+ column_filter::caller<20, T, D, BrdColConstant>,
+ column_filter::caller<21, T, D, BrdColConstant>,
+ column_filter::caller<22, T, D, BrdColConstant>,
+ column_filter::caller<23, T, D, BrdColConstant>,
+ column_filter::caller<24, T, D, BrdColConstant>,
+ column_filter::caller<25, T, D, BrdColConstant>,
+ column_filter::caller<26, T, D, BrdColConstant>,
+ column_filter::caller<27, T, D, BrdColConstant>,
+ column_filter::caller<28, T, D, BrdColConstant>,
+ column_filter::caller<29, T, D, BrdColConstant>,
+ column_filter::caller<30, T, D, BrdColConstant>,
+ column_filter::caller<31, T, D, BrdColConstant>,
+ column_filter::caller<32, T, D, BrdColConstant>
},
{
0,
},
{
0,
- column_filter::caller< 1, T, D, BrdColConstant>,
- column_filter::caller< 2, T, D, BrdColConstant>,
- column_filter::caller< 3, T, D, BrdColConstant>,
- column_filter::caller< 4, T, D, BrdColConstant>,
- column_filter::caller< 5, T, D, BrdColConstant>,
- column_filter::caller< 6, T, D, BrdColConstant>,
- column_filter::caller< 7, T, D, BrdColConstant>,
- column_filter::caller< 8, T, D, BrdColConstant>,
- column_filter::caller< 9, T, D, BrdColConstant>,
- column_filter::caller<10, T, D, BrdColConstant>,
- column_filter::caller<11, T, D, BrdColConstant>,
- column_filter::caller<12, T, D, BrdColConstant>,
- column_filter::caller<13, T, D, BrdColConstant>,
- column_filter::caller<14, T, D, BrdColConstant>,
- column_filter::caller<15, T, D, BrdColConstant>,
- column_filter::caller<16, T, D, BrdColConstant>,
- column_filter::caller<17, T, D, BrdColConstant>,
- column_filter::caller<18, T, D, BrdColConstant>,
- column_filter::caller<19, T, D, BrdColConstant>,
- column_filter::caller<20, T, D, BrdColConstant>,
- column_filter::caller<21, T, D, BrdColConstant>,
- column_filter::caller<22, T, D, BrdColConstant>,
- column_filter::caller<23, T, D, BrdColConstant>,
- column_filter::caller<24, T, D, BrdColConstant>,
- column_filter::caller<25, T, D, BrdColConstant>,
- column_filter::caller<26, T, D, BrdColConstant>,
- column_filter::caller<27, T, D, BrdColConstant>,
- column_filter::caller<28, T, D, BrdColConstant>,
- column_filter::caller<29, T, D, BrdColConstant>,
- column_filter::caller<30, T, D, BrdColConstant>,
- column_filter::caller<31, T, D, BrdColConstant>,
- column_filter::caller<32, T, D, BrdColConstant>
- },
- {
- 0,
column_filter::caller< 1, T, D, BrdColReflect>,
column_filter::caller< 2, T, D, BrdColReflect>,
column_filter::caller< 3, T, D, BrdColReflect>,
column_filter::caller<30, T, D, BrdColWrap>,
column_filter::caller<31, T, D, BrdColWrap>,
column_filter::caller<32, T, D, BrdColWrap>
+ },
+ {
+ 0,
+ column_filter::caller< 1, T, D, BrdColReflect101>,
+ column_filter::caller< 2, T, D, BrdColReflect101>,
+ column_filter::caller< 3, T, D, BrdColReflect101>,
+ column_filter::caller< 4, T, D, BrdColReflect101>,
+ column_filter::caller< 5, T, D, BrdColReflect101>,
+ column_filter::caller< 6, T, D, BrdColReflect101>,
+ column_filter::caller< 7, T, D, BrdColReflect101>,
+ column_filter::caller< 8, T, D, BrdColReflect101>,
+ column_filter::caller< 9, T, D, BrdColReflect101>,
+ column_filter::caller<10, T, D, BrdColReflect101>,
+ column_filter::caller<11, T, D, BrdColReflect101>,
+ column_filter::caller<12, T, D, BrdColReflect101>,
+ column_filter::caller<13, T, D, BrdColReflect101>,
+ column_filter::caller<14, T, D, BrdColReflect101>,
+ column_filter::caller<15, T, D, BrdColReflect101>,
+ column_filter::caller<16, T, D, BrdColReflect101>,
+ column_filter::caller<17, T, D, BrdColReflect101>,
+ column_filter::caller<18, T, D, BrdColReflect101>,
+ column_filter::caller<19, T, D, BrdColReflect101>,
+ column_filter::caller<20, T, D, BrdColReflect101>,
+ column_filter::caller<21, T, D, BrdColReflect101>,
+ column_filter::caller<22, T, D, BrdColReflect101>,
+ column_filter::caller<23, T, D, BrdColReflect101>,
+ column_filter::caller<24, T, D, BrdColReflect101>,
+ column_filter::caller<25, T, D, BrdColReflect101>,
+ column_filter::caller<26, T, D, BrdColReflect101>,
+ column_filter::caller<27, T, D, BrdColReflect101>,
+ column_filter::caller<28, T, D, BrdColReflect101>,
+ column_filter::caller<29, T, D, BrdColReflect101>,
+ column_filter::caller<30, T, D, BrdColReflect101>,
+ column_filter::caller<31, T, D, BrdColReflect101>,
+ column_filter::caller<32, T, D, BrdColReflect101>
}
};
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);
static const func_t funcs[] =
{
- Filter2DCaller<T, D, BrdReflect101>::call,
- Filter2DCaller<T, D, BrdReplicate>::call,
Filter2DCaller<T, D, BrdConstant>::call,
+ Filter2DCaller<T, D, BrdReplicate>::call,
Filter2DCaller<T, D, BrdReflect>::call,
- Filter2DCaller<T, D, BrdWrap>::call
+ Filter2DCaller<T, D, BrdWrap>::call,
+ Filter2DCaller<T, D, BrdReflect101>::call
};
if (stream == 0)
{
{
0,
- row_filter::caller< 1, T, D, BrdRowReflect101>,
- row_filter::caller< 2, T, D, BrdRowReflect101>,
- row_filter::caller< 3, T, D, BrdRowReflect101>,
- row_filter::caller< 4, T, D, BrdRowReflect101>,
- row_filter::caller< 5, T, D, BrdRowReflect101>,
- row_filter::caller< 6, T, D, BrdRowReflect101>,
- row_filter::caller< 7, T, D, BrdRowReflect101>,
- row_filter::caller< 8, T, D, BrdRowReflect101>,
- row_filter::caller< 9, T, D, BrdRowReflect101>,
- row_filter::caller<10, T, D, BrdRowReflect101>,
- row_filter::caller<11, T, D, BrdRowReflect101>,
- row_filter::caller<12, T, D, BrdRowReflect101>,
- row_filter::caller<13, T, D, BrdRowReflect101>,
- row_filter::caller<14, T, D, BrdRowReflect101>,
- row_filter::caller<15, T, D, BrdRowReflect101>,
- row_filter::caller<16, T, D, BrdRowReflect101>,
- row_filter::caller<17, T, D, BrdRowReflect101>,
- row_filter::caller<18, T, D, BrdRowReflect101>,
- row_filter::caller<19, T, D, BrdRowReflect101>,
- row_filter::caller<20, T, D, BrdRowReflect101>,
- row_filter::caller<21, T, D, BrdRowReflect101>,
- row_filter::caller<22, T, D, BrdRowReflect101>,
- row_filter::caller<23, T, D, BrdRowReflect101>,
- row_filter::caller<24, T, D, BrdRowReflect101>,
- row_filter::caller<25, T, D, BrdRowReflect101>,
- row_filter::caller<26, T, D, BrdRowReflect101>,
- row_filter::caller<27, T, D, BrdRowReflect101>,
- row_filter::caller<28, T, D, BrdRowReflect101>,
- row_filter::caller<29, T, D, BrdRowReflect101>,
- row_filter::caller<30, T, D, BrdRowReflect101>,
- row_filter::caller<31, T, D, BrdRowReflect101>,
- row_filter::caller<32, T, D, BrdRowReflect101>
+ row_filter::caller< 1, T, D, BrdRowConstant>,
+ row_filter::caller< 2, T, D, BrdRowConstant>,
+ row_filter::caller< 3, T, D, BrdRowConstant>,
+ row_filter::caller< 4, T, D, BrdRowConstant>,
+ row_filter::caller< 5, T, D, BrdRowConstant>,
+ row_filter::caller< 6, T, D, BrdRowConstant>,
+ row_filter::caller< 7, T, D, BrdRowConstant>,
+ row_filter::caller< 8, T, D, BrdRowConstant>,
+ row_filter::caller< 9, T, D, BrdRowConstant>,
+ row_filter::caller<10, T, D, BrdRowConstant>,
+ row_filter::caller<11, T, D, BrdRowConstant>,
+ row_filter::caller<12, T, D, BrdRowConstant>,
+ row_filter::caller<13, T, D, BrdRowConstant>,
+ row_filter::caller<14, T, D, BrdRowConstant>,
+ row_filter::caller<15, T, D, BrdRowConstant>,
+ row_filter::caller<16, T, D, BrdRowConstant>,
+ row_filter::caller<17, T, D, BrdRowConstant>,
+ row_filter::caller<18, T, D, BrdRowConstant>,
+ row_filter::caller<19, T, D, BrdRowConstant>,
+ row_filter::caller<20, T, D, BrdRowConstant>,
+ row_filter::caller<21, T, D, BrdRowConstant>,
+ row_filter::caller<22, T, D, BrdRowConstant>,
+ row_filter::caller<23, T, D, BrdRowConstant>,
+ row_filter::caller<24, T, D, BrdRowConstant>,
+ row_filter::caller<25, T, D, BrdRowConstant>,
+ row_filter::caller<26, T, D, BrdRowConstant>,
+ row_filter::caller<27, T, D, BrdRowConstant>,
+ row_filter::caller<28, T, D, BrdRowConstant>,
+ row_filter::caller<29, T, D, BrdRowConstant>,
+ row_filter::caller<30, T, D, BrdRowConstant>,
+ row_filter::caller<31, T, D, BrdRowConstant>,
+ row_filter::caller<32, T, D, BrdRowConstant>
},
{
0,
},
{
0,
- row_filter::caller< 1, T, D, BrdRowConstant>,
- row_filter::caller< 2, T, D, BrdRowConstant>,
- row_filter::caller< 3, T, D, BrdRowConstant>,
- row_filter::caller< 4, T, D, BrdRowConstant>,
- row_filter::caller< 5, T, D, BrdRowConstant>,
- row_filter::caller< 6, T, D, BrdRowConstant>,
- row_filter::caller< 7, T, D, BrdRowConstant>,
- row_filter::caller< 8, T, D, BrdRowConstant>,
- row_filter::caller< 9, T, D, BrdRowConstant>,
- row_filter::caller<10, T, D, BrdRowConstant>,
- row_filter::caller<11, T, D, BrdRowConstant>,
- row_filter::caller<12, T, D, BrdRowConstant>,
- row_filter::caller<13, T, D, BrdRowConstant>,
- row_filter::caller<14, T, D, BrdRowConstant>,
- row_filter::caller<15, T, D, BrdRowConstant>,
- row_filter::caller<16, T, D, BrdRowConstant>,
- row_filter::caller<17, T, D, BrdRowConstant>,
- row_filter::caller<18, T, D, BrdRowConstant>,
- row_filter::caller<19, T, D, BrdRowConstant>,
- row_filter::caller<20, T, D, BrdRowConstant>,
- row_filter::caller<21, T, D, BrdRowConstant>,
- row_filter::caller<22, T, D, BrdRowConstant>,
- row_filter::caller<23, T, D, BrdRowConstant>,
- row_filter::caller<24, T, D, BrdRowConstant>,
- row_filter::caller<25, T, D, BrdRowConstant>,
- row_filter::caller<26, T, D, BrdRowConstant>,
- row_filter::caller<27, T, D, BrdRowConstant>,
- row_filter::caller<28, T, D, BrdRowConstant>,
- row_filter::caller<29, T, D, BrdRowConstant>,
- row_filter::caller<30, T, D, BrdRowConstant>,
- row_filter::caller<31, T, D, BrdRowConstant>,
- row_filter::caller<32, T, D, BrdRowConstant>
- },
- {
- 0,
row_filter::caller< 1, T, D, BrdRowReflect>,
row_filter::caller< 2, T, D, BrdRowReflect>,
row_filter::caller< 3, T, D, BrdRowReflect>,
row_filter::caller<30, T, D, BrdRowWrap>,
row_filter::caller<31, T, D, BrdRowWrap>,
row_filter::caller<32, T, D, BrdRowWrap>
+ },
+ {
+ 0,
+ row_filter::caller< 1, T, D, BrdRowReflect101>,
+ row_filter::caller< 2, T, D, BrdRowReflect101>,
+ row_filter::caller< 3, T, D, BrdRowReflect101>,
+ row_filter::caller< 4, T, D, BrdRowReflect101>,
+ row_filter::caller< 5, T, D, BrdRowReflect101>,
+ row_filter::caller< 6, T, D, BrdRowReflect101>,
+ row_filter::caller< 7, T, D, BrdRowReflect101>,
+ row_filter::caller< 8, T, D, BrdRowReflect101>,
+ row_filter::caller< 9, T, D, BrdRowReflect101>,
+ row_filter::caller<10, T, D, BrdRowReflect101>,
+ row_filter::caller<11, T, D, BrdRowReflect101>,
+ row_filter::caller<12, T, D, BrdRowReflect101>,
+ row_filter::caller<13, T, D, BrdRowReflect101>,
+ row_filter::caller<14, T, D, BrdRowReflect101>,
+ row_filter::caller<15, T, D, BrdRowReflect101>,
+ row_filter::caller<16, T, D, BrdRowReflect101>,
+ row_filter::caller<17, T, D, BrdRowReflect101>,
+ row_filter::caller<18, T, D, BrdRowReflect101>,
+ row_filter::caller<19, T, D, BrdRowReflect101>,
+ row_filter::caller<20, T, D, BrdRowReflect101>,
+ row_filter::caller<21, T, D, BrdRowReflect101>,
+ row_filter::caller<22, T, D, BrdRowReflect101>,
+ row_filter::caller<23, T, D, BrdRowReflect101>,
+ row_filter::caller<24, T, D, BrdRowReflect101>,
+ row_filter::caller<25, T, D, BrdRowReflect101>,
+ row_filter::caller<26, T, D, BrdRowReflect101>,
+ row_filter::caller<27, T, D, BrdRowReflect101>,
+ row_filter::caller<28, T, D, BrdRowReflect101>,
+ row_filter::caller<29, T, D, BrdRowReflect101>,
+ row_filter::caller<30, T, D, BrdRowReflect101>,
+ row_filter::caller<31, T, D, BrdRowReflect101>,
+ row_filter::caller<32, T, D, BrdRowReflect101>
}
};
CV_Assert(ksize.width * ksize.height <= 16 * 16);
- int gpuBorderType;
- CV_Assert( tryConvertToGpuBorderType(brd_type, gpuBorderType) );
-
GpuMat gpu_krnl;
normalizeKernel(kernel, gpu_krnl, CV_32F);
break;
}
- return Ptr<BaseFilter_GPU>(new GpuFilter2D(ksize, anchor, func, gpu_krnl, gpuBorderType));
+ return Ptr<BaseFilter_GPU>(new GpuFilter2D(ksize, anchor, func, gpu_krnl, brd_type));
}
Ptr<FilterEngine_GPU> cv::gpu::createLinearFilter_GPU(int srcType, int dstType, const Mat& kernel, Point anchor, int borderType)
CV_Assert( borderType == BORDER_REFLECT101 || borderType == BORDER_REPLICATE || borderType == BORDER_CONSTANT || borderType == BORDER_REFLECT || borderType == BORDER_WRAP );
- int gpuBorderType;
- CV_Assert( tryConvertToGpuBorderType(borderType, gpuBorderType) );
-
const int sdepth = CV_MAT_DEPTH(srcType);
const int cn = CV_MAT_CN(srcType);
CV_Assert( sdepth <= CV_64F && cn <= 4 );
normalizeAnchor(anchor, ksize);
- return Ptr<BaseRowFilter_GPU>(new GpuLinearRowFilter(ksize, anchor, gpu_row_krnl, func, gpuBorderType));
+ return Ptr<BaseRowFilter_GPU>(new GpuLinearRowFilter(ksize, anchor, gpu_row_krnl, func, borderType));
}
namespace
CV_Assert( borderType == BORDER_REFLECT101 || borderType == BORDER_REPLICATE || borderType == BORDER_CONSTANT || borderType == BORDER_REFLECT || borderType == BORDER_WRAP );
- int gpuBorderType;
- CV_Assert( tryConvertToGpuBorderType(borderType, gpuBorderType) );
-
const int ddepth = CV_MAT_DEPTH(dstType);
const int cn = CV_MAT_CN(dstType);
CV_Assert( ddepth <= CV_64F && cn <= 4 );
normalizeAnchor(anchor, ksize);
- return Ptr<BaseColumnFilter_GPU>(new GpuLinearColumnFilter(ksize, anchor, gpu_col_krnl, func, gpuBorderType));
+ return Ptr<BaseColumnFilter_GPU>(new GpuLinearColumnFilter(ksize, anchor, gpu_col_krnl, func, borderType));
}
Ptr<FilterEngine_GPU> cv::gpu::createSeparableLinearFilter_GPU(int srcType, int dstType, const Mat& rowKernel, const Mat& columnKernel,
CV_Assert(borderMode == BORDER_REFLECT101 || borderMode == BORDER_REPLICATE || borderMode == BORDER_CONSTANT || borderMode == BORDER_REFLECT || borderMode == BORDER_WRAP);
- int gpuBorderType;
- CV_Assert(tryConvertToGpuBorderType(borderMode, gpuBorderType));
-
dst.create(src.size(), src.type());
- func(src, dst, kernel_size, sigma_spatial, sigma_color, gpuBorderType, StreamAccessor::getStream(s));
+ func(src, dst, kernel_size, sigma_spatial, sigma_color, borderMode, StreamAccessor::getStream(s));
}
#endif
static caller_t funcs[] =
{
- bilateral_caller<T, BrdReflect101>,
- bilateral_caller<T, BrdReplicate>,
bilateral_caller<T, BrdConstant>,
+ bilateral_caller<T, BrdReplicate>,
bilateral_caller<T, BrdReflect>,
bilateral_caller<T, BrdWrap>,
+ bilateral_caller<T, BrdReflect101>
};
funcs[borderMode](src, dst, kernel_size, gauss_spatial_coeff, gauss_color_coeff, stream);
}
switch (border_type)
{
- case BORDER_REFLECT101_GPU:
+ case BORDER_REFLECT101:
cornerHarris_kernel<<<grid, block, 0, stream>>>(block_size, k, dst, BrdRowReflect101<void>(Dx.cols), BrdColReflect101<void>(Dx.rows));
break;
- case BORDER_REFLECT_GPU:
+ case BORDER_REFLECT:
cornerHarris_kernel<<<grid, block, 0, stream>>>(block_size, k, dst, BrdRowReflect<void>(Dx.cols), BrdColReflect<void>(Dx.rows));
break;
- case BORDER_REPLICATE_GPU:
+ case BORDER_REPLICATE:
cornerHarris_kernel<<<grid, block, 0, stream>>>(block_size, k, dst);
break;
}
switch (border_type)
{
- case BORDER_REFLECT101_GPU:
+ case BORDER_REFLECT101:
cornerMinEigenVal_kernel<<<grid, block, 0, stream>>>(block_size, dst, BrdRowReflect101<void>(Dx.cols), BrdColReflect101<void>(Dx.rows));
break;
- case BORDER_REFLECT_GPU:
+ case BORDER_REFLECT:
cornerMinEigenVal_kernel<<<grid, block, 0, stream>>>(block_size, dst, BrdRowReflect<void>(Dx.cols), BrdColReflect<void>(Dx.rows));
break;
- case BORDER_REPLICATE_GPU:
+ case BORDER_REPLICATE:
cornerMinEigenVal_kernel<<<grid, block, 0, stream>>>(block_size, dst);
break;
}
CV_Assert(borderType == cv::BORDER_REFLECT101 || borderType == cv::BORDER_REPLICATE || borderType == cv::BORDER_REFLECT);
- int gpuBorderType;
- CV_Assert(tryConvertToGpuBorderType(borderType, gpuBorderType));
-
extractCovData(src, Dx, Dy, buf, blockSize, ksize, borderType, stream);
dst.create(src.size(), CV_32F);
- cornerHarris_gpu(blockSize, static_cast<float>(k), Dx, Dy, dst, gpuBorderType, StreamAccessor::getStream(stream));
+ cornerHarris_gpu(blockSize, static_cast<float>(k), Dx, Dy, dst, borderType, StreamAccessor::getStream(stream));
}
void cv::gpu::cornerMinEigenVal(const GpuMat& src, GpuMat& dst, int blockSize, int ksize, int borderType)
CV_Assert(borderType == cv::BORDER_REFLECT101 || borderType == cv::BORDER_REPLICATE || borderType == cv::BORDER_REFLECT);
- int gpuBorderType;
- CV_Assert(tryConvertToGpuBorderType(borderType, gpuBorderType));
-
extractCovData(src, Dx, Dy, buf, blockSize, ksize, borderType, stream);
dst.create(src.size(), CV_32F);
- cornerMinEigenVal_gpu(blockSize, Dx, Dy, dst, gpuBorderType, StreamAccessor::getStream(stream));
+ cornerMinEigenVal_gpu(blockSize, Dx, Dy, dst, borderType, StreamAccessor::getStream(stream));
}
static const caller_t callers[] =
{
- gaussianBlurCaller<BrdReflect101<float> >,
+ 0 /*gaussianBlurCaller<BrdConstant<float> >*/,
gaussianBlurCaller<BrdReplicate<float> >,
+ 0 /*gaussianBlurCaller<BrdReflect<float> >*/,
+ 0 /*gaussianBlurCaller<BrdWrap<float> >*/,
+ gaussianBlurCaller<BrdReflect101<float> >
};
callers[borderMode](src, ksizeHalf, dst, stream);
static const caller_t callers[] =
{
- gaussianBlur5Caller<BrdReflect101<float>,256>,
+ 0 /*gaussianBlur5Caller<BrdConstant<float>,256>*/,
gaussianBlur5Caller<BrdReplicate<float>,256>,
+ 0 /*gaussianBlur5Caller<BrdReflect<float>,256>*/,
+ 0 /*gaussianBlur5Caller<BrdWrap<float>,256>*/,
+ gaussianBlur5Caller<BrdReflect101<float>,256>
};
callers[borderMode](src, ksizeHalf, dst, stream);
static const caller_t callers[] =
{
- gaussianBlur5Caller<BrdReflect101<float>,128>,
+ 0 /*gaussianBlur5Caller<BrdConstant<float>,128>*/,
gaussianBlur5Caller<BrdReplicate<float>,128>,
+ 0 /*gaussianBlur5Caller<BrdReflect<float>,128>*/,
+ 0 /*gaussianBlur5Caller<BrdWrap<float>,128>*/,
+ gaussianBlur5Caller<BrdReflect101<float>,128>
};
callers[borderMode](src, ksizeHalf, dst, stream);
{
if (deviceSupports(FEATURE_SET_COMPUTE_12))
cudev::optflow_farneback::gaussianBlur5Gpu(
- M, blockSize/2, bufM, BORDER_REPLICATE_GPU, S(streams[0]));
+ M, blockSize/2, bufM, BORDER_REPLICATE, S(streams[0]));
else
cudev::optflow_farneback::gaussianBlur5Gpu_CC11(
- M, blockSize/2, bufM, BORDER_REPLICATE_GPU, S(streams[0]));
+ M, blockSize/2, bufM, BORDER_REPLICATE, S(streams[0]));
swap(M, bufM);
cudev::optflow_farneback::updateFlowGpu(M, flowx, flowy, S(streams[0]));
for (int i = 0; i < 2; i++)
{
cudev::optflow_farneback::gaussianBlurGpu(
- frames_[i], smoothSize/2, blurredFrame[i], BORDER_REFLECT101_GPU, S(streams[i]));
+ frames_[i], smoothSize/2, blurredFrame[i], BORDER_REFLECT101, S(streams[i]));
#if ENABLE_GPU_RESIZE
resize(blurredFrame[i], pyrLevel[i], Size(width, height), INTER_LINEAR, streams[i]);
#else
static const caller_t callers[3][5] =
{
{
- RemapDispatcher<PointFilter, BrdReflect101, T>::call,
- RemapDispatcher<PointFilter, BrdReplicate, T>::call,
RemapDispatcher<PointFilter, BrdConstant, T>::call,
+ RemapDispatcher<PointFilter, BrdReplicate, T>::call,
RemapDispatcher<PointFilter, BrdReflect, T>::call,
- RemapDispatcher<PointFilter, BrdWrap, T>::call
+ RemapDispatcher<PointFilter, BrdWrap, T>::call,
+ RemapDispatcher<PointFilter, BrdReflect101, T>::call
},
{
- RemapDispatcher<LinearFilter, BrdReflect101, T>::call,
- RemapDispatcher<LinearFilter, BrdReplicate, T>::call,
RemapDispatcher<LinearFilter, BrdConstant, T>::call,
+ RemapDispatcher<LinearFilter, BrdReplicate, T>::call,
RemapDispatcher<LinearFilter, BrdReflect, T>::call,
- RemapDispatcher<LinearFilter, BrdWrap, T>::call
+ RemapDispatcher<LinearFilter, BrdWrap, T>::call,
+ RemapDispatcher<LinearFilter, BrdReflect101, T>::call
},
{
- RemapDispatcher<CubicFilter, BrdReflect101, T>::call,
- RemapDispatcher<CubicFilter, BrdReplicate, T>::call,
RemapDispatcher<CubicFilter, BrdConstant, T>::call,
+ RemapDispatcher<CubicFilter, BrdReplicate, T>::call,
RemapDispatcher<CubicFilter, BrdReflect, T>::call,
- RemapDispatcher<CubicFilter, BrdWrap, T>::call
+ RemapDispatcher<CubicFilter, BrdWrap, T>::call,
+ RemapDispatcher<CubicFilter, BrdReflect101, T>::call
}
};
static const func_t funcs[3][5] =
{
{
- WarpDispatcher<Transform, PointFilter, BrdReflect101, T>::call,
- WarpDispatcher<Transform, PointFilter, BrdReplicate, T>::call,
WarpDispatcher<Transform, PointFilter, BrdConstant, T>::call,
+ WarpDispatcher<Transform, PointFilter, BrdReplicate, T>::call,
WarpDispatcher<Transform, PointFilter, BrdReflect, T>::call,
- WarpDispatcher<Transform, PointFilter, BrdWrap, T>::call
+ WarpDispatcher<Transform, PointFilter, BrdWrap, T>::call,
+ WarpDispatcher<Transform, PointFilter, BrdReflect101, T>::call
},
{
- WarpDispatcher<Transform, LinearFilter, BrdReflect101, T>::call,
- WarpDispatcher<Transform, LinearFilter, BrdReplicate, T>::call,
WarpDispatcher<Transform, LinearFilter, BrdConstant, T>::call,
+ WarpDispatcher<Transform, LinearFilter, BrdReplicate, T>::call,
WarpDispatcher<Transform, LinearFilter, BrdReflect, T>::call,
- WarpDispatcher<Transform, LinearFilter, BrdWrap, T>::call
+ WarpDispatcher<Transform, LinearFilter, BrdWrap, T>::call,
+ WarpDispatcher<Transform, LinearFilter, BrdReflect101, T>::call
},
{
- WarpDispatcher<Transform, CubicFilter, BrdReflect101, T>::call,
- WarpDispatcher<Transform, CubicFilter, BrdReplicate, T>::call,
WarpDispatcher<Transform, CubicFilter, BrdConstant, T>::call,
+ WarpDispatcher<Transform, CubicFilter, BrdReplicate, T>::call,
WarpDispatcher<Transform, CubicFilter, BrdReflect, T>::call,
- WarpDispatcher<Transform, CubicFilter, BrdWrap, T>::call
+ WarpDispatcher<Transform, CubicFilter, BrdWrap, T>::call,
+ WarpDispatcher<Transform, CubicFilter, BrdReflect101, T>::call
}
};
const func_t func = funcs[src.depth()][src.channels() - 1];
CV_Assert(func != 0);
- int gpuBorderType;
- CV_Assert(tryConvertToGpuBorderType(borderMode, gpuBorderType));
-
dst.create(xmap.size(), src.type());
Scalar_<float> borderValueFloat;
src.locateROI(wholeSize, ofs);
func(src, PtrStepSzb(wholeSize.height, wholeSize.width, src.datastart, src.step), ofs.x, ofs.y, xmap, ymap,
- dst, interpolation, gpuBorderType, borderValueFloat.val, StreamAccessor::getStream(stream), deviceSupports(FEATURE_SET_COMPUTE_20));
+ dst, interpolation, borderMode, borderValueFloat.val, StreamAccessor::getStream(stream), deviceSupports(FEATURE_SET_COMPUTE_20));
}
#endif // HAVE_CUDA
const func_t func = funcs[src.depth()][src.channels() - 1];
CV_Assert(func != 0);
- int gpuBorderType;
- CV_Assert(tryConvertToGpuBorderType(borderMode, gpuBorderType));
-
float coeffs[2 * 3];
Mat coeffsMat(2, 3, CV_32F, (void*)coeffs);
borderValueFloat = borderValue;
func(src, PtrStepSzb(wholeSize.height, wholeSize.width, src.datastart, src.step), ofs.x, ofs.y, coeffs,
- dst, interpolation, gpuBorderType, borderValueFloat.val, StreamAccessor::getStream(s), deviceSupports(FEATURE_SET_COMPUTE_20));
+ dst, interpolation, borderMode, borderValueFloat.val, StreamAccessor::getStream(s), deviceSupports(FEATURE_SET_COMPUTE_20));
}
}
const func_t func = funcs[src.depth()][src.channels() - 1];
CV_Assert(func != 0);
- int gpuBorderType;
- CV_Assert(tryConvertToGpuBorderType(borderMode, gpuBorderType));
-
float coeffs[3 * 3];
Mat coeffsMat(3, 3, CV_32F, (void*)coeffs);
borderValueFloat = borderValue;
func(src, PtrStepSzb(wholeSize.height, wholeSize.width, src.datastart, src.step), ofs.x, ofs.y, coeffs,
- dst, interpolation, gpuBorderType, borderValueFloat.val, StreamAccessor::getStream(s), deviceSupports(FEATURE_SET_COMPUTE_20));
+ dst, interpolation, borderMode, borderValueFloat.val, StreamAccessor::getStream(s), deviceSupports(FEATURE_SET_COMPUTE_20));
}
}
static func_t funcs[] =
{
- nlm_caller<T, BrdReflect101>,
- nlm_caller<T, BrdReplicate>,
nlm_caller<T, BrdConstant>,
+ nlm_caller<T, BrdReplicate>,
nlm_caller<T, BrdReflect>,
nlm_caller<T, BrdWrap>,
+ nlm_caller<T, BrdReflect101>
};
funcs[borderMode](src, dst, search_radius, block_radius, h, stream);
}
int b = borderMode;
CV_Assert(b == BORDER_REFLECT101 || b == BORDER_REPLICATE || b == BORDER_CONSTANT || b == BORDER_REFLECT || b == BORDER_WRAP);
- int gpuBorderType;
- CV_Assert(tryConvertToGpuBorderType(borderMode, gpuBorderType));
-
dst.create(src.size(), src.type());
- func(src, dst, search_window/2, block_window/2, h, gpuBorderType, StreamAccessor::getStream(s));
+ func(src, dst, search_window/2, block_window/2, h, borderMode, StreamAccessor::getStream(s));
}
namespace cv { namespace gpu { namespace cudev