diff --git a/modules/core/src/convert.cpp b/modules/core/src/convert.cpp index 05c1a6e402..6259a7ada2 100644 --- a/modules/core/src/convert.cpp +++ b/modules/core/src/convert.cpp @@ -1357,43 +1357,65 @@ void cv::LUT( InputArray _src, InputArray _lut, OutputArray _dst ) func(ptrs[0], lut.data, ptrs[1], len, cn, lutcn); } +namespace cv { + +static bool ocl_normalize( InputArray _src, OutputArray _dst, InputArray _mask, int rtype, + double scale, double shift ) +{ + UMat src = _src.getUMat(), dst = _dst.getUMat(); + + if( _mask.empty() ) + src.convertTo( dst, rtype, scale, shift ); + else + { + UMat temp; + src.convertTo( temp, rtype, scale, shift ); + temp.copyTo( dst, _mask ); + } + + return true; +} + +} void cv::normalize( InputArray _src, OutputArray _dst, double a, double b, int norm_type, int rtype, InputArray _mask ) { - Mat src = _src.getMat(), mask = _mask.getMat(); - double scale = 1, shift = 0; if( norm_type == CV_MINMAX ) { double smin = 0, smax = 0; double dmin = MIN( a, b ), dmax = MAX( a, b ); - minMaxLoc( _src, &smin, &smax, 0, 0, mask ); + minMaxLoc( _src, &smin, &smax, 0, 0, _mask ); scale = (dmax - dmin)*(smax - smin > DBL_EPSILON ? 1./(smax - smin) : 0); shift = dmin - smin*scale; } else if( norm_type == CV_L2 || norm_type == CV_L1 || norm_type == CV_C ) { - scale = norm( src, norm_type, mask ); + scale = norm( _src, norm_type, _mask ); scale = scale > DBL_EPSILON ? a/scale : 0.; shift = 0; } else CV_Error( CV_StsBadArg, "Unknown/unsupported norm type" ); + int type = _src.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type); if( rtype < 0 ) - rtype = _dst.fixedType() ? _dst.depth() : src.depth(); + rtype = _dst.fixedType() ? _dst.depth() : depth; + _dst.createSameSize(_src, CV_MAKETYPE(rtype, cn)); - _dst.create(src.dims, src.size, CV_MAKETYPE(rtype, src.channels())); - Mat dst = _dst.getMat(); + if (ocl::useOpenCL() && _dst.isUMat() && + ocl_normalize(_src, _dst, _mask, rtype, scale, shift)) + return; - if( !mask.data ) + Mat src = _src.getMat(), dst = _dst.getMat(); + if( _mask.empty() ) src.convertTo( dst, rtype, scale, shift ); else { Mat temp; src.convertTo( temp, rtype, scale, shift ); - temp.copyTo( dst, mask ); + temp.copyTo( dst, _mask ); } } diff --git a/modules/core/src/opencl/copyset.cl b/modules/core/src/opencl/copyset.cl index 8fb5a00cfe..05cde8ee01 100644 --- a/modules/core/src/opencl/copyset.cl +++ b/modules/core/src/opencl/copyset.cl @@ -41,6 +41,52 @@ // //M*/ +#ifdef COPY_TO_MASK + +#define DEFINE_DATA \ + int src_index = mad24(y, src_step, x*(int)sizeof(T)*scn + src_offset); \ + int dst_index = mad24(y, dst_step, x*(int)sizeof(T)*scn + dst_offset); \ + \ + __global const T * src = (__global const T *)(srcptr + src_index); \ + __global T * dst = (__global T *)(dstptr + dst_index) + +__kernel void copyToMask(__global const uchar * srcptr, int src_step, int src_offset, + __global const uchar * maskptr, int mask_step, int mask_offset, + __global uchar * dstptr, int dst_step, int dst_offset, + int dst_rows, int dst_cols) +{ + int x = get_global_id(0); + int y = get_global_id(1); + + if (x < dst_cols && y < dst_rows) + { + int mask_index = mad24(y, mask_step, x * mcn + mask_offset); + __global const uchar * mask = (__global const uchar *)(maskptr + mask_index); + +#if mcn == 1 + if (mask[0]) + { + DEFINE_DATA; + + #pragma unroll + for (int c = 0; c < scn; ++c) + dst[c] = src[c]; + } +#elif scn == mcn + DEFINE_DATA; + + #pragma unroll + for (int c = 0; c < scn; ++c) + if (mask[c]) + dst[c] = src[c]; +#else +#error "(mcn == 1 || mcn == scn) should be true" +#endif + } +} + +#else + __kernel void setMask(__global const uchar* mask, int maskstep, int maskoffset, __global uchar* dstptr, int dststep, int dstoffset, int rows, int cols, dstT value ) @@ -71,3 +117,5 @@ __kernel void set(__global uchar* dstptr, int dststep, int dstoffset, *(__global dstT*)(dstptr + dst_index) = value; } } + +#endif diff --git a/modules/core/src/umatrix.cpp b/modules/core/src/umatrix.cpp index 0b6137488a..95e203be9a 100644 --- a/modules/core/src/umatrix.cpp +++ b/modules/core/src/umatrix.cpp @@ -661,6 +661,45 @@ void UMat::copyTo(OutputArray _dst) const } } +void UMat::copyTo(OutputArray _dst, InputArray _mask) const +{ + if( _mask.empty() ) + { + copyTo(_dst); + return; + } + + int cn = channels(), mtype = _mask.type(), mdepth = CV_MAT_DEPTH(mtype), mcn = CV_MAT_CN(mtype); + CV_Assert( mdepth == CV_8U && (mcn == 1 || mcn == cn) ); + + if (ocl::useOpenCL() && _dst.isUMat() && dims <= 2) + { + UMatData * prevu = _dst.getUMat().u; + _dst.create( dims, size, type() ); + + UMat dst = _dst.getUMat(); + + if( prevu != dst.u ) // do not leave dst uninitialized + dst = Scalar(0); + + ocl::Kernel k("copyToMask", ocl::core::copyset_oclsrc, + format("-D COPY_TO_MASK -D T=%s -D scn=%d -D mcn=%d", + ocl::memopTypeToStr(depth()), cn, mcn)); + if (!k.empty()) + { + k.args(ocl::KernelArg::ReadOnlyNoSize(*this), ocl::KernelArg::ReadOnlyNoSize(_mask.getUMat()), + ocl::KernelArg::WriteOnly(dst)); + + size_t globalsize[2] = { cols, rows }; + if (k.run(2, globalsize, NULL, false)) + return; + } + } + + Mat src = getMat(ACCESS_READ); + src.copyTo(_dst, _mask); +} + void UMat::convertTo(OutputArray _dst, int _type, double alpha, double beta) const { bool noScale = std::fabs(alpha - 1) < DBL_EPSILON && std::fabs(beta) < DBL_EPSILON; diff --git a/modules/core/test/ocl/test_arithm.cpp b/modules/core/test/ocl/test_arithm.cpp index 0c37c71164..58edceccd2 100644 --- a/modules/core/test/ocl/test_arithm.cpp +++ b/modules/core/test/ocl/test_arithm.cpp @@ -1219,6 +1219,28 @@ OCL_TEST_P(Sqrt, Mat) } } +//////////////////////////////// Normalize //////////////////////////////////////////////// + +typedef ArithmTestBase Normalize; + +OCL_TEST_P(Normalize, Mat) +{ + static int modes[] = { CV_MINMAX, CV_L2, CV_L1, CV_C }; + + for (int j = 0; j < test_loop_times; j++) + { + generateTestData(); + + for (int i = 0, size = sizeof(modes) / sizeof(modes[0]); i < size; ++i) + { + OCL_OFF(cv::normalize(src1_roi, dst1_roi, 10, 110, modes[i], src1_roi.type(), mask_roi)); + OCL_ON(cv::normalize(usrc1_roi, udst1_roi, 10, 110, modes[i], src1_roi.type(), umask_roi)); + + Near(1); + } + } +} + //////////////////////////////////////// Instantiation ///////////////////////////////////////// OCL_INSTANTIATE_TEST_CASE_P(Arithm, Lut, Combine(::testing::Values(CV_8U, CV_8S), OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool(), Bool())); @@ -1253,6 +1275,8 @@ OCL_INSTANTIATE_TEST_CASE_P(Arithm, MinMaxIdx, Combine(OCL_ALL_DEPTHS, OCL_ALL_C OCL_INSTANTIATE_TEST_CASE_P(Arithm, MinMaxIdx_Mask, Combine(OCL_ALL_DEPTHS, ::testing::Values(Channels(1)), Bool())); OCL_INSTANTIATE_TEST_CASE_P(Arithm, Norm, Combine(OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool())); OCL_INSTANTIATE_TEST_CASE_P(Arithm, Sqrt, Combine(::testing::Values(CV_32F, CV_64F), OCL_ALL_CHANNELS, Bool())); +OCL_INSTANTIATE_TEST_CASE_P(Arithm, Normalize, Combine(OCL_ALL_DEPTHS, Values(Channels(1)), Bool())); + } } // namespace cvtest::ocl diff --git a/modules/core/test/ocl/test_matrix_operation.cpp b/modules/core/test/ocl/test_matrix_operation.cpp index aabbb3f6c3..77c5dad956 100644 --- a/modules/core/test/ocl/test_matrix_operation.cpp +++ b/modules/core/test/ocl/test_matrix_operation.cpp @@ -54,7 +54,7 @@ namespace ocl { ////////////////////////////////converto///////////////////////////////////////////////// -PARAM_TEST_CASE(MatrixTestBase, MatDepth, MatDepth, Channels, bool) +PARAM_TEST_CASE(ConvertTo, MatDepth, MatDepth, Channels, bool) { int src_depth, cn, dstType; bool use_roi; @@ -85,8 +85,6 @@ PARAM_TEST_CASE(MatrixTestBase, MatDepth, MatDepth, Channels, bool) } }; -typedef MatrixTestBase ConvertTo; - OCL_TEST_P(ConvertTo, Accuracy) { for (int j = 0; j < test_loop_times; j++) @@ -103,7 +101,51 @@ OCL_TEST_P(ConvertTo, Accuracy) } } -typedef MatrixTestBase CopyTo; +//////////////////////////////// CopyTo ///////////////////////////////////////////////// + +PARAM_TEST_CASE(CopyTo, MatDepth, Channels, bool, bool) +{ + int depth, cn; + bool use_roi, use_mask; + + TEST_DECLARE_INPUT_PARAMETER(src) + TEST_DECLARE_INPUT_PARAMETER(mask) + TEST_DECLARE_OUTPUT_PARAMETER(dst) + + virtual void SetUp() + { + depth = GET_PARAM(0); + cn = GET_PARAM(1); + use_roi = GET_PARAM(2); + use_mask = GET_PARAM(3); + } + + void generateTestData() + { + const int type = CV_MAKE_TYPE(depth, cn); + + Size roiSize = randomSize(1, MAX_VALUE); + Border srcBorder = randomBorder(0, use_roi ? MAX_VALUE : 0); + randomSubMat(src, src_roi, roiSize, srcBorder, type, -MAX_VALUE, MAX_VALUE); + + if (use_mask) + { + Border maskBorder = randomBorder(0, use_roi ? MAX_VALUE : 0); + int mask_cn = randomDouble(0.0, 2.0) > 1.0 ? cn : 1; + randomSubMat(mask, mask_roi, roiSize, maskBorder, CV_8UC(mask_cn), 0, 2); + cv::threshold(mask, mask, 0.5, 255., CV_8UC1); + } + + Border dstBorder = randomBorder(0, use_roi ? MAX_VALUE : 0); + randomSubMat(dst, dst_roi, roiSize, dstBorder, type, 5, 16); + + UMAT_UPLOAD_INPUT_PARAMETER(src) + if (use_mask) + UMAT_UPLOAD_INPUT_PARAMETER(mask) + UMAT_UPLOAD_OUTPUT_PARAMETER(dst) + } +}; + OCL_TEST_P(CopyTo, Accuracy) { @@ -111,8 +153,16 @@ OCL_TEST_P(CopyTo, Accuracy) { generateTestData(); - OCL_OFF(src_roi.copyTo(dst_roi)); - OCL_ON(usrc_roi.copyTo(udst_roi)); + if (use_mask) + { + OCL_OFF(src_roi.copyTo(dst_roi, mask_roi)); + OCL_ON(usrc_roi.copyTo(udst_roi, umask_roi)); + } + else + { + OCL_OFF(src_roi.copyTo(dst_roi)); + OCL_ON(usrc_roi.copyTo(udst_roi)); + } OCL_EXPECT_MATS_NEAR(dst, 0); } @@ -122,7 +172,7 @@ OCL_INSTANTIATE_TEST_CASE_P(MatrixOperation, ConvertTo, Combine( OCL_ALL_DEPTHS, OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool())); OCL_INSTANTIATE_TEST_CASE_P(MatrixOperation, CopyTo, Combine( - OCL_ALL_DEPTHS, Values((MatDepth)0), OCL_ALL_CHANNELS, Bool())); + OCL_ALL_DEPTHS, OCL_ALL_CHANNELS, Bool(), Bool())); } } // namespace cvtest::ocl