}\r
\r
GPU_SANITY_CHECK(d_vertex);\r
- GPU_SANITY_CHECK(d_colors)\r
+ GPU_SANITY_CHECK(d_colors);\r
}\r
else\r
{\r
\r
__constant__ float c_kernel[MAX_KERNEL_SIZE];\r
\r
- void loadKernel(const float kernel[], int ksize)\r
+ void loadKernel(const float* kernel, int ksize, cudaStream_t stream)\r
{\r
- cudaSafeCall( cudaMemcpyToSymbol(c_kernel, kernel, ksize * sizeof(float)) );\r
+ if (stream == 0)\r
+ cudaSafeCall( cudaMemcpyToSymbol(c_kernel, kernel, ksize * sizeof(float), 0, cudaMemcpyDeviceToDevice) );\r
+ else\r
+ cudaSafeCall( cudaMemcpyToSymbolAsync(c_kernel, kernel, ksize * sizeof(float), 0, cudaMemcpyDeviceToDevice, stream) );\r
}\r
\r
template <int KSIZE, typename T, typename D, typename B>\r
}\r
\r
template <typename T, typename D>\r
- void linearColumnFilter_gpu(PtrStepSzb src, PtrStepSzb dst, const float kernel[], int ksize, int anchor, int brd_type, int cc, cudaStream_t stream)\r
+ void linearColumnFilter_gpu(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream)\r
{\r
typedef void (*caller_t)(PtrStepSz<T> src, PtrStepSz<D> dst, int anchor, int cc, cudaStream_t stream);\r
\r
}\r
};\r
\r
- loadKernel(kernel, ksize);\r
+ loadKernel(kernel, ksize, stream);\r
\r
callers[brd_type][ksize]((PtrStepSz<T>)src, (PtrStepSz<D>)dst, anchor, cc, stream);\r
}\r
\r
- template void linearColumnFilter_gpu<float , uchar >(PtrStepSzb src, PtrStepSzb dst, const float kernel[], int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);\r
- template void linearColumnFilter_gpu<float4, uchar4>(PtrStepSzb src, PtrStepSzb dst, const float kernel[], int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);\r
- template void linearColumnFilter_gpu<float3, short3>(PtrStepSzb src, PtrStepSzb dst, const float kernel[], int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);\r
- template void linearColumnFilter_gpu<float , int >(PtrStepSzb src, PtrStepSzb dst, const float kernel[], int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);\r
- template void linearColumnFilter_gpu<float , float >(PtrStepSzb src, PtrStepSzb dst, const float kernel[], int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);\r
+ template void linearColumnFilter_gpu<float , uchar >(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);\r
+ template void linearColumnFilter_gpu<float4, uchar4>(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);\r
+ template void linearColumnFilter_gpu<float3, short3>(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);\r
+ template void linearColumnFilter_gpu<float , int >(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);\r
+ template void linearColumnFilter_gpu<float , float >(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);\r
} // namespace column_filter\r
}}} // namespace cv { namespace gpu { namespace device\r
\r
\r
-#endif /* CUDA_DISABLER */
\ No newline at end of file
+#endif /* CUDA_DISABLER */\r
Filter2DCaller<T, D, BrdWrap>::call\r
};\r
\r
- cudaSafeCall(cudaMemcpyToSymbol(c_filter2DKernel, kernel, kWidth * kHeight * sizeof(float), 0, cudaMemcpyDeviceToDevice) );\r
+ if (stream == 0)\r
+ cudaSafeCall( cudaMemcpyToSymbol(c_filter2DKernel, kernel, kWidth * kHeight * sizeof(float), 0, cudaMemcpyDeviceToDevice) );\r
+ else\r
+ cudaSafeCall( cudaMemcpyToSymbolAsync(c_filter2DKernel, kernel, kWidth * kHeight * sizeof(float), 0, cudaMemcpyDeviceToDevice, stream) );\r
\r
funcs[borderMode](static_cast< PtrStepSz<T> >(srcWhole), ofsX, ofsY, static_cast< PtrStepSz<D> >(dst), kWidth, kHeight, anchorX, anchorY, borderValue, stream);\r
}\r
}}} // namespace cv { namespace gpu { namespace device {\r
\r
\r
-#endif /* CUDA_DISABLER */
\ No newline at end of file
+#endif /* CUDA_DISABLER */\r
\r
__constant__ float c_kernel[MAX_KERNEL_SIZE];\r
\r
- void loadKernel(const float kernel[], int ksize)\r
+ void loadKernel(const float* kernel, int ksize, cudaStream_t stream)\r
{\r
- cudaSafeCall( cudaMemcpyToSymbol(c_kernel, kernel, ksize * sizeof(float)) );\r
+ if (stream == 0)\r
+ cudaSafeCall( cudaMemcpyToSymbol(c_kernel, kernel, ksize * sizeof(float), 0, cudaMemcpyDeviceToDevice) );\r
+ else\r
+ cudaSafeCall( cudaMemcpyToSymbolAsync(c_kernel, kernel, ksize * sizeof(float), 0, cudaMemcpyDeviceToDevice, stream) );\r
}\r
\r
template <int KSIZE, typename T, typename D, typename B>\r
}\r
\r
template <typename T, typename D>\r
- void linearRowFilter_gpu(PtrStepSzb src, PtrStepSzb dst, const float kernel[], int ksize, int anchor, int brd_type, int cc, cudaStream_t stream)\r
+ void linearRowFilter_gpu(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream)\r
{\r
typedef void (*caller_t)(PtrStepSz<T> src, PtrStepSz<D> dst, int anchor, int cc, cudaStream_t stream);\r
\r
}\r
};\r
\r
- loadKernel(kernel, ksize);\r
+ loadKernel(kernel, ksize, stream);\r
\r
callers[brd_type][ksize]((PtrStepSz<T>)src, (PtrStepSz<D>)dst, anchor, cc, stream);\r
}\r
\r
- template void linearRowFilter_gpu<uchar , float >(PtrStepSzb src, PtrStepSzb dst, const float kernel[], int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);\r
- template void linearRowFilter_gpu<uchar4, float4>(PtrStepSzb src, PtrStepSzb dst, const float kernel[], int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);\r
- template void linearRowFilter_gpu<short3, float3>(PtrStepSzb src, PtrStepSzb dst, const float kernel[], int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);\r
- template void linearRowFilter_gpu<int , float >(PtrStepSzb src, PtrStepSzb dst, const float kernel[], int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);\r
- template void linearRowFilter_gpu<float , float >(PtrStepSzb src, PtrStepSzb dst, const float kernel[], int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);\r
+ template void linearRowFilter_gpu<uchar , float >(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);\r
+ template void linearRowFilter_gpu<uchar4, float4>(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);\r
+ template void linearRowFilter_gpu<short3, float3>(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);\r
+ template void linearRowFilter_gpu<int , float >(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);\r
+ template void linearRowFilter_gpu<float , float >(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);\r
} // namespace row_filter\r
}}} // namespace cv { namespace gpu { namespace device\r
\r
\r
-#endif /* CUDA_DISABLER */
\ No newline at end of file
+#endif /* CUDA_DISABLER */\r
namespace row_filter\r
{\r
template <typename T, typename D>\r
- void linearRowFilter_gpu(PtrStepSzb src, PtrStepSzb dst, const float kernel[], int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);\r
+ void linearRowFilter_gpu(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);\r
}\r
\r
namespace column_filter\r
{\r
template <typename T, typename D>\r
- void linearColumnFilter_gpu(PtrStepSzb src, PtrStepSzb dst, const float kernel[], int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);\r
+ void linearColumnFilter_gpu(PtrStepSzb src, PtrStepSzb dst, const float* kernel, int ksize, int anchor, int brd_type, int cc, cudaStream_t stream);\r
}\r
}}}\r
\r
\r
struct GpuLinearRowFilter : public BaseRowFilter_GPU\r
{\r
- GpuLinearRowFilter(int ksize_, int anchor_, const Mat& kernel_, gpuFilter1D_t func_, int brd_type_) :\r
+ GpuLinearRowFilter(int ksize_, int anchor_, const GpuMat& kernel_, gpuFilter1D_t func_, int brd_type_) :\r
BaseRowFilter_GPU(ksize_, anchor_), kernel(kernel_), func(func_), brd_type(brd_type_) {}\r
\r
virtual void operator()(const GpuMat& src, GpuMat& dst, Stream& s = Stream::Null())\r
func(src, dst, kernel.ptr<float>(), ksize, anchor, brd_type, cc, StreamAccessor::getStream(s));\r
}\r
\r
- Mat kernel;\r
+ GpuMat kernel;\r
gpuFilter1D_t func;\r
int brd_type;\r
};\r
\r
CV_Assert(CV_MAT_DEPTH(bufType) == CV_32F && CV_MAT_CN(srcType) == CV_MAT_CN(bufType));\r
\r
- Mat temp(rowKernel.size(), CV_32FC1);\r
- rowKernel.convertTo(temp, CV_32FC1);\r
- Mat cont_krnl = temp.reshape(1, 1);\r
+ GpuMat gpu_row_krnl;\r
+ normalizeKernel(rowKernel, gpu_row_krnl, CV_32F);\r
\r
- int ksize = cont_krnl.cols;\r
+ int ksize = gpu_row_krnl.cols;\r
\r
CV_Assert(ksize > 0 && ksize <= 32);\r
\r
break;\r
}\r
\r
- return Ptr<BaseRowFilter_GPU>(new GpuLinearRowFilter(ksize, anchor, cont_krnl, func, gpuBorderType));\r
+ return Ptr<BaseRowFilter_GPU>(new GpuLinearRowFilter(ksize, anchor, gpu_row_krnl, func, gpuBorderType));\r
}\r
\r
namespace\r
\r
struct GpuLinearColumnFilter : public BaseColumnFilter_GPU\r
{\r
- GpuLinearColumnFilter(int ksize_, int anchor_, const Mat& kernel_, gpuFilter1D_t func_, int brd_type_) :\r
+ GpuLinearColumnFilter(int ksize_, int anchor_, const GpuMat& kernel_, gpuFilter1D_t func_, int brd_type_) :\r
BaseColumnFilter_GPU(ksize_, anchor_), kernel(kernel_), func(func_), brd_type(brd_type_) {}\r
\r
virtual void operator()(const GpuMat& src, GpuMat& dst, Stream& s = Stream::Null())\r
func(src, dst, kernel.ptr<float>(), ksize, anchor, brd_type, cc, StreamAccessor::getStream(s));\r
}\r
\r
- Mat kernel;\r
+ GpuMat kernel;\r
gpuFilter1D_t func;\r
int brd_type;\r
};\r
\r
CV_Assert(CV_MAT_DEPTH(bufType) == CV_32F && CV_MAT_CN(dstType) == CV_MAT_CN(bufType));\r
\r
- Mat temp(columnKernel.size(), CV_32FC1);\r
- columnKernel.convertTo(temp, CV_32FC1);\r
- Mat cont_krnl = temp.reshape(1, 1);\r
+ GpuMat gpu_col_krnl;\r
+ normalizeKernel(columnKernel, gpu_col_krnl, CV_32F);\r
\r
- int ksize = cont_krnl.cols;\r
+ int ksize = gpu_col_krnl.cols;\r
\r
CV_Assert(ksize > 0 && ksize <= 32);\r
\r
break;\r
}\r
\r
- return Ptr<BaseColumnFilter_GPU>(new GpuLinearColumnFilter(ksize, anchor, cont_krnl, func, gpuBorderType));\r
+ return Ptr<BaseColumnFilter_GPU>(new GpuLinearColumnFilter(ksize, anchor, gpu_col_krnl, func, gpuBorderType));\r
}\r
\r
Ptr<FilterEngine_GPU> cv::gpu::createSeparableLinearFilter_GPU(int srcType, int dstType, const Mat& rowKernel, const Mat& columnKernel,\r