diff --git a/modules/core/include/opencv2/core.hpp b/modules/core/include/opencv2/core.hpp index bd5de32d8d..cfb995b1da 100644 --- a/modules/core/include/opencv2/core.hpp +++ b/modules/core/include/opencv2/core.hpp @@ -556,6 +556,8 @@ The format of half precision floating point is defined in IEEE 754-2008. @param src input array. @param dst output array. + +@deprecated Use Mat::convertTo with CV_16F instead. */ CV_EXPORTS_W void convertFp16(InputArray src, OutputArray dst); diff --git a/modules/core/include/opencv2/core/ocl.hpp b/modules/core/include/opencv2/core/ocl.hpp index 0000343fa7..891fd678b7 100644 --- a/modules/core/include/opencv2/core/ocl.hpp +++ b/modules/core/include/opencv2/core/ocl.hpp @@ -127,6 +127,11 @@ public: CV_WRAP int singleFPConfig() const; CV_WRAP int halfFPConfig() const; + /// true if 'cl_khr_fp64' extension is available + CV_WRAP bool hasFP64() const; + /// true if 'cl_khr_fp16' extension is available + CV_WRAP bool hasFP16() const; + CV_WRAP bool endianLittle() const; CV_WRAP bool errorCorrectionSupport() const; diff --git a/modules/core/include/opencv2/core/opencl/opencl_info.hpp b/modules/core/include/opencv2/core/opencl/opencl_info.hpp index 0f0de893ca..845efba9fc 100644 --- a/modules/core/include/opencv2/core/opencl/opencl_info.hpp +++ b/modules/core/include/opencv2/core/opencl/opencl_info.hpp @@ -141,13 +141,13 @@ static void dumpOpenCLInformation() DUMP_MESSAGE_STDOUT(" Max memory allocation size = " << maxMemAllocSizeStr); DUMP_CONFIG_PROPERTY("cv_ocl_current_maxMemAllocSize", device.maxMemAllocSize()); - const char* doubleSupportStr = device.doubleFPConfig() > 0 ? "Yes" : "No"; + const char* doubleSupportStr = device.hasFP64() ? "Yes" : "No"; DUMP_MESSAGE_STDOUT(" Double support = " << doubleSupportStr); - DUMP_CONFIG_PROPERTY("cv_ocl_current_haveDoubleSupport", device.doubleFPConfig() > 0); + DUMP_CONFIG_PROPERTY("cv_ocl_current_haveDoubleSupport", device.hasFP64()); - const char* halfSupportStr = device.halfFPConfig() > 0 ? "Yes" : "No"; + const char* halfSupportStr = device.hasFP16() ? "Yes" : "No"; DUMP_MESSAGE_STDOUT(" Half support = " << halfSupportStr); - DUMP_CONFIG_PROPERTY("cv_ocl_current_haveHalfSupport", device.halfFPConfig() > 0); + DUMP_CONFIG_PROPERTY("cv_ocl_current_haveHalfSupport", device.hasFP16()); const char* isUnifiedMemoryStr = device.hostUnifiedMemory() ? "Yes" : "No"; DUMP_MESSAGE_STDOUT(" Host unified memory = " << isUnifiedMemoryStr); diff --git a/modules/core/perf/opencl/perf_matop.cpp b/modules/core/perf/opencl/perf_matop.cpp index b763a98e2a..2fd5b177a6 100644 --- a/modules/core/perf/opencl/perf_matop.cpp +++ b/modules/core/perf/opencl/perf_matop.cpp @@ -80,6 +80,187 @@ OCL_PERF_TEST_P(ConvertToFixture, ConvertTo, SANITY_CHECK(dst); } + +//#define RUN_CONVERTFP16 +static Size convertFP16_srcSize(4000, 4000); + +OCL_PERF_TEST(Core, ConvertFP32FP16MatMat) +{ + const Size srcSize = convertFP16_srcSize; + const int type = CV_32F; + const int dtype = CV_16F; + + checkDeviceMaxMemoryAllocSize(srcSize, type); + checkDeviceMaxMemoryAllocSize(srcSize, dtype); + + Mat src(srcSize, type); + Mat dst(srcSize, dtype); + declare.in(src, WARMUP_RNG).out(dst); + +#ifdef RUN_CONVERTFP16 + OCL_TEST_CYCLE() convertFp16(src, dst); +#else + OCL_TEST_CYCLE() src.convertTo(dst, dtype); +#endif + + SANITY_CHECK_NOTHING(); +} + +OCL_PERF_TEST(Core, ConvertFP32FP16MatUMat) +{ + const Size srcSize = convertFP16_srcSize; + const int type = CV_32F; + const int dtype = CV_16F; + + checkDeviceMaxMemoryAllocSize(srcSize, type); + checkDeviceMaxMemoryAllocSize(srcSize, dtype); + + Mat src(srcSize, type); + UMat dst(srcSize, dtype); + declare.in(src, WARMUP_RNG).out(dst); + +#ifdef RUN_CONVERTFP16 + OCL_TEST_CYCLE() convertFp16(src, dst); +#else + OCL_TEST_CYCLE() src.convertTo(dst, dtype); +#endif + + SANITY_CHECK_NOTHING(); +} + +OCL_PERF_TEST(Core, ConvertFP32FP16UMatMat) +{ + const Size srcSize = convertFP16_srcSize; + const int type = CV_32F; + const int dtype = CV_16F; + + checkDeviceMaxMemoryAllocSize(srcSize, type); + checkDeviceMaxMemoryAllocSize(srcSize, dtype); + + UMat src(srcSize, type); + Mat dst(srcSize, dtype); + declare.in(src, WARMUP_RNG).out(dst); + +#ifdef RUN_CONVERTFP16 + OCL_TEST_CYCLE() convertFp16(src, dst); +#else + OCL_TEST_CYCLE() src.convertTo(dst, dtype); +#endif + + SANITY_CHECK_NOTHING(); +} + +OCL_PERF_TEST(Core, ConvertFP32FP16UMatUMat) +{ + const Size srcSize = convertFP16_srcSize; + const int type = CV_32F; + const int dtype = CV_16F; + + checkDeviceMaxMemoryAllocSize(srcSize, type); + checkDeviceMaxMemoryAllocSize(srcSize, dtype); + + UMat src(srcSize, type); + UMat dst(srcSize, dtype); + declare.in(src, WARMUP_RNG).out(dst); + +#ifdef RUN_CONVERTFP16 + OCL_TEST_CYCLE() convertFp16(src, dst); +#else + OCL_TEST_CYCLE() src.convertTo(dst, dtype); +#endif + + SANITY_CHECK_NOTHING(); +} + +OCL_PERF_TEST(Core, ConvertFP16FP32MatMat) +{ + const Size srcSize = convertFP16_srcSize; + const int type = CV_16F; + const int dtype = CV_32F; + + checkDeviceMaxMemoryAllocSize(srcSize, type); + checkDeviceMaxMemoryAllocSize(srcSize, dtype); + + Mat src(srcSize, type); + Mat dst(srcSize, dtype); + declare.in(src, WARMUP_RNG).out(dst); + +#ifdef RUN_CONVERTFP16 + OCL_TEST_CYCLE() convertFp16(src, dst); +#else + OCL_TEST_CYCLE() src.convertTo(dst, dtype); +#endif + + SANITY_CHECK_NOTHING(); +} + +OCL_PERF_TEST(Core, ConvertFP16FP32MatUMat) +{ + const Size srcSize = convertFP16_srcSize; + const int type = CV_16F; + const int dtype = CV_32F; + + checkDeviceMaxMemoryAllocSize(srcSize, type); + checkDeviceMaxMemoryAllocSize(srcSize, dtype); + + Mat src(srcSize, type); + UMat dst(srcSize, dtype); + declare.in(src, WARMUP_RNG).out(dst); + +#ifdef RUN_CONVERTFP16 + OCL_TEST_CYCLE() convertFp16(src, dst); +#else + OCL_TEST_CYCLE() src.convertTo(dst, dtype); +#endif + + SANITY_CHECK_NOTHING(); +} + +OCL_PERF_TEST(Core, ConvertFP16FP32UMatMat) +{ + const Size srcSize = convertFP16_srcSize; + const int type = CV_16F; + const int dtype = CV_32F; + + checkDeviceMaxMemoryAllocSize(srcSize, type); + checkDeviceMaxMemoryAllocSize(srcSize, dtype); + + UMat src(srcSize, type); + Mat dst(srcSize, dtype); + declare.in(src, WARMUP_RNG).out(dst); + +#ifdef RUN_CONVERTFP16 + OCL_TEST_CYCLE() convertFp16(src, dst); +#else + OCL_TEST_CYCLE() src.convertTo(dst, dtype); +#endif + + SANITY_CHECK_NOTHING(); +} + +OCL_PERF_TEST(Core, ConvertFP16FP32UMatUMat) +{ + const Size srcSize = convertFP16_srcSize; + const int type = CV_16F; + const int dtype = CV_32F; + + checkDeviceMaxMemoryAllocSize(srcSize, type); + checkDeviceMaxMemoryAllocSize(srcSize, dtype); + + UMat src(srcSize, type); + UMat dst(srcSize, dtype); + declare.in(src, WARMUP_RNG).out(dst); + +#ifdef RUN_CONVERTFP16 + OCL_TEST_CYCLE() convertFp16(src, dst); +#else + OCL_TEST_CYCLE() src.convertTo(dst, dtype); +#endif + + SANITY_CHECK_NOTHING(); +} + + ///////////// CopyTo //////////////////////// typedef Size_MatType CopyToFixture; diff --git a/modules/core/src/convert.dispatch.cpp b/modules/core/src/convert.dispatch.cpp index 345b4624cb..58ca43187a 100644 --- a/modules/core/src/convert.dispatch.cpp +++ b/modules/core/src/convert.dispatch.cpp @@ -169,52 +169,130 @@ static bool ocl_convertFp16( InputArray _src, OutputArray _dst, int sdepth, int size_t globalsize[2] = { (size_t)src.cols * cn / kercn, ((size_t)src.rows + rowsPerWI - 1) / rowsPerWI }; return k.run(2, globalsize, NULL, false); } -#endif -void Mat::convertTo(OutputArray _dst, int _type, double alpha, double beta) const +static bool ocl_convertTo(InputArray src_, OutputArray dst_, int ddepth, bool noScale, double alpha, double beta) { CV_INSTRUMENT_REGION(); - if( empty() ) + CV_Assert(ddepth >= 0); + + int stype = src_.type(); + int sdepth = CV_MAT_DEPTH(stype); + int cn = CV_MAT_CN(stype); + + int dtype = CV_MAKETYPE(ddepth, cn); + + int wdepth = (sdepth == CV_64F) ? CV_64F : CV_32F; + + bool needDouble = sdepth == CV_64F || ddepth == CV_64F; + bool doubleCheck = true; + if (needDouble) { - _dst.release(); - return; + doubleCheck = ocl::Device::getDefault().hasFP64(); + } + bool halfCheck = true; + bool needHalf = sdepth == CV_16F || ddepth == CV_16F; + if (needHalf) + { + halfCheck = ocl::Device::getDefault().hasFP16(); } - bool noScale = fabs(alpha-1) < DBL_EPSILON && fabs(beta) < DBL_EPSILON; + if (!doubleCheck) + return false; + if (!halfCheck) + return false; - if( _type < 0 ) - _type = _dst.fixedType() ? _dst.type() : type(); + const int rowsPerWI = 4; + + char cvt[2][50]; + ocl::Kernel k("convertTo", ocl::core::convert_oclsrc, + format("-D srcT=%s -D WT=%s -D dstT=%s -D convertToWT=%s -D convertToDT=%s -D rowsPerWI=%d%s%s%s", + ocl::typeToStr(sdepth), ocl::typeToStr(wdepth), ocl::typeToStr(ddepth), + ocl::convertTypeStr(sdepth, wdepth, 1, cvt[0], sizeof(cvt[0])), + ocl::convertTypeStr(wdepth, ddepth, 1, cvt[1], sizeof(cvt[1])), + rowsPerWI, + needDouble ? " -D DOUBLE_SUPPORT" : "", + needHalf ? " -D HALF_SUPPORT" : "", + noScale ? " -D NO_SCALE" : "" + ) + ); + + if (k.empty()) + return false; + + UMat src = src_.getUMat(); + dst_.createSameSize(src_, dtype); + UMat dst = dst_.getUMat(); + + float alphaf = (float)alpha, betaf = (float)beta; + + if (noScale) + k.args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::WriteOnly(dst, cn)); + else if (wdepth == CV_32F) + k.args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::WriteOnly(dst, cn), alphaf, betaf); else - _type = CV_MAKETYPE(CV_MAT_DEPTH(_type), channels()); + k.args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::WriteOnly(dst, cn), alpha, beta); - int sdepth = depth(), ddepth = CV_MAT_DEPTH(_type); - if( sdepth == ddepth && noScale ) + size_t globalsize[2] = { + (size_t)dst.cols * cn, + divUp((size_t)dst.rows, rowsPerWI) + }; + if (!k.run(2, globalsize, NULL, false)) + return false; + + CV_IMPL_ADD(CV_IMPL_OCL); + return true; +} +#endif + +void Mat::convertTo(OutputArray dst, int type_, double alpha, double beta) const +{ + CV_INSTRUMENT_REGION(); + + if (empty()) { - copyTo(_dst); + dst.release(); return; } + int stype = type(); + int sdepth = CV_MAT_DEPTH(stype); + + int ddepth = sdepth; + if (type_ >= 0) + ddepth = CV_MAT_DEPTH(type_); + else + ddepth = dst.fixedType() ? dst.depth() : sdepth; + + bool noScale = std::fabs(alpha - 1) < DBL_EPSILON && std::fabs(beta) < DBL_EPSILON; + if (sdepth == ddepth && noScale) + { + copyTo(dst); + return; + } + + CV_OCL_RUN(dims <= 2 && dst.isUMat(), + ocl_convertTo(*this, dst, ddepth, noScale, alpha, beta)) + + int cn = channels(); + int dtype = CV_MAKETYPE(ddepth, cn); + Mat src = *this; - if( dims <= 2 ) - _dst.create( size(), _type ); - else - _dst.create( dims, size, _type ); - Mat dst = _dst.getMat(); + dst.create(dims, size, dtype); + Mat dstMat = dst.getMat(); BinaryFunc func = noScale ? getConvertFunc(sdepth, ddepth) : getConvertScaleFunc(sdepth, ddepth); double scale[] = {alpha, beta}; - int cn = channels(); CV_Assert( func != 0 ); if( dims <= 2 ) { - Size sz = getContinuousSize2D(src, dst, cn); - func( src.data, src.step, 0, 0, dst.data, dst.step, sz, scale ); + Size sz = getContinuousSize2D(src, dstMat, cn); + func(src.data, src.step, 0, 0, dstMat.data, dstMat.step, sz, scale); } else { - const Mat* arrays[] = {&src, &dst, 0}; + const Mat* arrays[] = {&src, &dstMat, 0}; uchar* ptrs[2] = {}; NAryMatIterator it(arrays, ptrs); Size sz((int)(it.size*cn), 1); @@ -224,6 +302,44 @@ void Mat::convertTo(OutputArray _dst, int _type, double alpha, double beta) cons } } +void UMat::convertTo(OutputArray dst, int type_, double alpha, double beta) const +{ + CV_INSTRUMENT_REGION(); + + if (empty()) + { + dst.release(); + return; + } + +#ifdef HAVE_OPENCL + int stype = type(); + int sdepth = CV_MAT_DEPTH(stype); + + int ddepth = sdepth; + if (type_ >= 0) + ddepth = CV_MAT_DEPTH(type_); + else + ddepth = dst.fixedType() ? dst.depth() : sdepth; + + bool noScale = std::fabs(alpha - 1) < DBL_EPSILON && std::fabs(beta) < DBL_EPSILON; + if (sdepth == ddepth && noScale) + { + copyTo(dst); + return; + } + + CV_OCL_RUN(dims <= 2, + ocl_convertTo(*this, dst, ddepth, noScale, alpha, beta)) +#endif // HAVE_OPENCL + + UMat src = *this; // Fake reference to itself. + // Resolves issue 8693 in case of src == dst. + Mat m = getMat(ACCESS_READ); + m.convertTo(dst, type_, alpha, beta); + (void)src; +} + //================================================================================================== void convertFp16(InputArray _src, OutputArray _dst) diff --git a/modules/core/src/ocl.cpp b/modules/core/src/ocl.cpp index f93a7be3f1..5eac178316 100644 --- a/modules/core/src/ocl.cpp +++ b/modules/core/src/ocl.cpp @@ -1604,6 +1604,9 @@ struct Device::Impl pos = pos2 + 1; } + khr_fp64_support_ = isExtensionSupported("cl_khr_fp64"); + khr_fp16_support_ = isExtensionSupported("cl_khr_fp16"); + intelSubgroupsSupport_ = isExtensionSupported("cl_intel_subgroups"); vendorName_ = getStrProp(CL_DEVICE_VENDOR); @@ -1692,7 +1695,9 @@ struct Device::Impl String version_; std::string extensions_; int doubleFPConfig_; + bool khr_fp64_support_; int halfFPConfig_; + bool khr_fp16_support_; bool hostUnifiedMemory_; int maxComputeUnits_; size_t maxWorkGroupSize_; @@ -1844,6 +1849,11 @@ int Device::singleFPConfig() const int Device::halfFPConfig() const { return p ? p->halfFPConfig_ : 0; } +bool Device::hasFP64() const +{ return p ? p->khr_fp64_support_ : false; } +bool Device::hasFP16() const +{ return p ? p->khr_fp16_support_ : false; } + bool Device::endianLittle() const { return p ? p->getBoolProp(CL_DEVICE_ENDIAN_LITTLE) : false; } diff --git a/modules/core/src/ocl_disabled.impl.hpp b/modules/core/src/ocl_disabled.impl.hpp index a217979a1e..fab5351025 100644 --- a/modules/core/src/ocl_disabled.impl.hpp +++ b/modules/core/src/ocl_disabled.impl.hpp @@ -67,6 +67,9 @@ int Device::doubleFPConfig() const { OCL_NOT_AVAILABLE(); } int Device::singleFPConfig() const { OCL_NOT_AVAILABLE(); } int Device::halfFPConfig() const { OCL_NOT_AVAILABLE(); } +bool Device::hasFP64() const { OCL_NOT_AVAILABLE(); } +bool Device::hasFP16() const { OCL_NOT_AVAILABLE(); } + bool Device::endianLittle() const { OCL_NOT_AVAILABLE(); } bool Device::errorCorrectionSupport() const { OCL_NOT_AVAILABLE(); } diff --git a/modules/core/src/opencl/convert.cl b/modules/core/src/opencl/convert.cl index e869d6d743..1f58f63f07 100644 --- a/modules/core/src/opencl/convert.cl +++ b/modules/core/src/opencl/convert.cl @@ -49,14 +49,21 @@ #endif #endif +#ifdef HALF_SUPPORT +#ifdef cl_khr_fp16 +#pragma OPENCL EXTENSION cl_khr_fp16:enable +#endif +#endif + + #define noconvert __kernel void convertTo(__global const uchar * srcptr, int src_step, int src_offset, - __global uchar * dstptr, int dst_step, int dst_offset, int dst_rows, int dst_cols, + __global uchar * dstptr, int dst_step, int dst_offset, int dst_rows, int dst_cols #ifndef NO_SCALE - WT alpha, WT beta, + , WT alpha, WT beta #endif - int rowsPerWI) +) { int x = get_global_id(0); int y0 = get_global_id(1) * rowsPerWI; diff --git a/modules/core/src/umatrix.cpp b/modules/core/src/umatrix.cpp index 02b37026f4..5d3dbc5f0b 100644 --- a/modules/core/src/umatrix.cpp +++ b/modules/core/src/umatrix.cpp @@ -1233,70 +1233,10 @@ void UMat::copyTo(OutputArray _dst, InputArray _mask) const src.copyTo(_dst, _mask); } -void UMat::convertTo(OutputArray _dst, int _type, double alpha, double beta) const -{ - CV_INSTRUMENT_REGION(); - bool noScale = std::fabs(alpha - 1) < DBL_EPSILON && std::fabs(beta) < DBL_EPSILON; - int stype = type(), cn = CV_MAT_CN(stype); - - if( _type < 0 ) - _type = _dst.fixedType() ? _dst.type() : stype; - else - _type = CV_MAKETYPE(CV_MAT_DEPTH(_type), cn); - - int sdepth = CV_MAT_DEPTH(stype), ddepth = CV_MAT_DEPTH(_type); - if( sdepth == ddepth && noScale ) - { - copyTo(_dst); - return; - } -#ifdef HAVE_OPENCL - bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0; - bool needDouble = sdepth == CV_64F || ddepth == CV_64F; - if( dims <= 2 && cn && _dst.isUMat() && ocl::useOpenCL() && - ((needDouble && doubleSupport) || !needDouble) ) - { - int wdepth = std::max(CV_32F, sdepth), rowsPerWI = 4; - - char cvt[2][50]; - ocl::Kernel k("convertTo", ocl::core::convert_oclsrc, - format("-D srcT=%s -D WT=%s -D dstT=%s -D convertToWT=%s -D convertToDT=%s%s%s", - ocl::typeToStr(sdepth), ocl::typeToStr(wdepth), ocl::typeToStr(ddepth), - ocl::convertTypeStr(sdepth, wdepth, 1, cvt[0], sizeof(cvt[0])), - ocl::convertTypeStr(wdepth, ddepth, 1, cvt[1], sizeof(cvt[1])), - doubleSupport ? " -D DOUBLE_SUPPORT" : "", noScale ? " -D NO_SCALE" : "")); - if (!k.empty()) - { - UMat src = *this; - _dst.create( size(), _type ); - UMat dst = _dst.getUMat(); - - float alphaf = (float)alpha, betaf = (float)beta; - ocl::KernelArg srcarg = ocl::KernelArg::ReadOnlyNoSize(src), - dstarg = ocl::KernelArg::WriteOnly(dst, cn); - - if (noScale) - k.args(srcarg, dstarg, rowsPerWI); - else if (wdepth == CV_32F) - k.args(srcarg, dstarg, alphaf, betaf, rowsPerWI); - else - k.args(srcarg, dstarg, alpha, beta, rowsPerWI); - - size_t globalsize[2] = { (size_t)dst.cols * cn, ((size_t)dst.rows + rowsPerWI - 1) / rowsPerWI }; - if (k.run(2, globalsize, NULL, false)) - { - CV_IMPL_ADD(CV_IMPL_OCL); - return; - } - } - } -#endif - UMat src = *this; // Fake reference to itself. - // Resolves issue 8693 in case of src == dst. - Mat m = getMat(ACCESS_READ); - m.convertTo(_dst, _type, alpha, beta); -} +// +// void UMat::convertTo moved to convert.dispatch.cpp +// UMat& UMat::setTo(InputArray _value, InputArray _mask) {