summary | shortlog | log | commit | commitdiff | tree
raw | patch | inline | side by side (parent: b9a2c79)
raw | patch | inline | side by side (parent: b9a2c79)
author | Djordje Senicic <d-senicic1@ti.com> | |
Tue, 13 Sep 2016 12:18:50 +0000 (08:18 -0400) | ||
committer | Djordje Senicic <d-senicic1@ti.com> | |
Tue, 13 Sep 2016 12:18:50 +0000 (08:18 -0400) |
modules/imgproc/src/opencl/gauss.cl | patch | blob | history | |
modules/imgproc/src/smooth.cpp | patch | blob | history |
index 52b7ede92a34056f3c5e2e8ce27e1f877f4f8116..3c21d45a076e1881209bd36ede3a0d42fb25e35b 100644 (file)
unsigned int mask_r1_1 = mask_r1_0 << 8U;
if (!evIN) { printf("Failed to alloc edmaIN handle.\n"); return; }
-
+#ifdef TIDSP_OPENCL_VERBOSE
clk_start = __clock();
+#endif
rows >>= 1;
dest_ptr = (uchar *)dstptr;
@@ -144,7 +145,9 @@ __attribute__((reqd_work_group_size(1,1,1))) __kernel void tidsp_gaussian(__glob
EdmaMgr_wait(evIN);
EdmaMgr_free(evIN);
+#ifdef TIDSP_OPENCL_VERBOSE
clk_end = __clock();
printf ("TIDSP gauss clockdiff=%d\n", clk_end - clk_start);
+#endif
}
/********************************************************************************************/
index 62502373e014a9ae4517b76c1bd49b68a1bbcb41..a3dc85fdabddb2c1d129848e849563b4455d42f1 100644 (file)
size_t localsize[2] = { 1, 1 };
size_t globalsize[2] = { 2, 1 };
int type = _src.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type);
-
-// if ( !((depth == CV_8U || depth == CV_16U || depth == CV_16S || depth == CV_32F) && cn <= 4 && (m == 3 || m == 5)) )
-// return false;
+ int dtype = _dst.type(), ddepth = CV_MAT_DEPTH(dtype);
Size imgSize = _src.size();
bool useOptimized = (1 == cn) && (imgSize.width % 8 == 0) && ((size_t)imgSize.width >= 16);
- if( (ksize.width != 3) && (ksize.height != 3) ) useOptimized = false;
+ useOptimized &&= (ksize.width == 3) && (ksize.height == 3);
+ useOptimized &&= (depth == CV_8U) && (ddepth == CV_8U);
cv::String kname = format( "tidsp_gaussian" ) ;
cv::String kdefs = format("-D T=%s -D T1=%s -D cn=%d", ocl::typeToStr(type), ocl::typeToStr(depth), cn) ;
ocl::Kernel k(kname.c_str(), ocl::imgproc::gauss_oclsrc, kdefs.c_str() );
-printf ("\nTIDSP_GAUSSIAN:%d optimized:%d\n", !k.empty(), useOptimized);
if (!k.empty() && useOptimized)
{
UMat src = _src.getUMat();
#endif
CV_IPP_RUN(true, ipp_GaussianBlur( _src, _dst, ksize, sigma1, sigma2, borderType));
-printf ("\nDJDBG GaussianBlur (sigma1=%lf sigma2=%lf, borderType=%d)!\n", sigma1, sigma2, borderType); fflush(stdout);
Mat kx, ky;
createGaussianKernels(kx, ky, type, ksize, sigma1, sigma2);
sepFilter2D(_src, _dst, CV_MAT_DEPTH(type), kx, ky, Point(-1,-1), 0, borderType );
@@ -3162,13 +3159,13 @@ static bool ocl_bilateralFilter_8u(InputArray _src, OutputArray _dst, int d,
Mat mspace_ofs(1, d * d, CV_32SC1, space_ofs);
UMat ucolor_weight, uspace_weight, uspace_ofs;
-printf ("\nDJDBG bilateralFilter!\n");
mspace_weight.copyTo(uspace_weight);
mspace_ofs.copyTo(uspace_ofs);
k.args(ocl::KernelArg::ReadOnlyNoSize(temp), ocl::KernelArg::WriteOnly(dst),
ocl::KernelArg::PtrReadOnly(uspace_weight),
ocl::KernelArg::PtrReadOnly(uspace_ofs));
+
size_t globalsize[2] = { (size_t)dst.cols / sizeDiv, (size_t)dst.rows };
return k.run(2, globalsize, NULL, false);
}