From 878dec65c89c566f38987ea877c868d5dd0d046f Mon Sep 17 00:00:00 2001 From: vbystricky Date: Mon, 28 Jul 2014 15:48:02 +0400 Subject: [PATCH] Optimize OpenCL version of morfology and box filters for small filter kernels --- modules/imgproc/src/morph.cpp | 176 ++++++++++++++++- .../opencl/{boxFilterSmall.cl => filterSmall.cl} | 212 ++++++++++++++++----- modules/imgproc/src/smooth.cpp | 4 +- modules/imgproc/test/ocl/test_filters.cpp | 70 ++++++- 4 files changed, 399 insertions(+), 63 deletions(-) rename modules/imgproc/src/opencl/{boxFilterSmall.cl => filterSmall.cl} (71%) diff --git a/modules/imgproc/src/morph.cpp b/modules/imgproc/src/morph.cpp index 4f696b4..328d03d 100644 --- a/modules/imgproc/src/morph.cpp +++ b/modules/imgproc/src/morph.cpp @@ -1339,20 +1339,188 @@ static bool IPPMorphOp(int op, InputArray _src, OutputArray _dst, #ifdef HAVE_OPENCL +#define ROUNDUP(sz, n) ((sz) + (n) - 1 - (((sz) + (n) - 1) % (n))) + +static bool ocl_morphSmall( InputArray _src, OutputArray _dst, InputArray _kernel, Point anchor, int borderType, + int op, int actual_op = -1, InputArray _extraMat = noArray()) +{ + const ocl::Device & dev = ocl::Device::getDefault(); + int type = _src.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type), esz = CV_ELEM_SIZE(type); + bool doubleSupport = dev.doubleFPConfig() > 0; + + if (cn > 4 || (!doubleSupport && depth == CV_64F) || + _src.offset() % esz != 0 || _src.step() % esz != 0) + return false; + + Size ksize = _kernel.size(); + if (anchor.x < 0) + anchor.x = ksize.width / 2; + if (anchor.y < 0) + anchor.y = ksize.height / 2; + + Size size = _src.size(), wholeSize; + bool isolated = (borderType & BORDER_ISOLATED) != 0; + borderType &= ~BORDER_ISOLATED; + int wdepth = depth, wtype = type; + if (depth == CV_8U) + { + wdepth = CV_32S; + wtype = CV_MAKETYPE(wdepth, cn); + } + char cvt[2][40]; + + bool haveExtraMat = !_extraMat.empty(); + CV_Assert(actual_op <= 3 || haveExtraMat); + + const char * const borderMap[] = { "BORDER_CONSTANT", "BORDER_REPLICATE", + "BORDER_REFLECT", 0, "BORDER_REFLECT_101" }; + size_t globalsize[2] = { size.width, size.height }; + + UMat src = _src.getUMat(); + if (!isolated) + { + Point ofs; + src.locateROI(wholeSize, ofs); + } + + int h = isolated ? size.height : wholeSize.height; + int w = isolated ? size.width : wholeSize.width; + if (w < ksize.width || h < ksize.height) + return false; + + // Figure out what vector size to use for loading the pixels. + int pxLoadNumPixels = cn != 1 || size.width % 4 ? 1 : 4; + int pxLoadVecSize = cn * pxLoadNumPixels; + + // Figure out how many pixels per work item to compute in X and Y + // directions. Too many and we run out of registers. + int pxPerWorkItemX = 1, pxPerWorkItemY = 1; + if (cn <= 2 && ksize.width <= 4 && ksize.height <= 4) + { + pxPerWorkItemX = size.width % 8 ? size.width % 4 ? size.width % 2 ? 1 : 2 : 4 : 8; + pxPerWorkItemY = size.height % 2 ? 1 : 2; + } + else if (cn < 4 || (ksize.width <= 4 && ksize.height <= 4)) + { + pxPerWorkItemX = size.width % 2 ? 1 : 2; + pxPerWorkItemY = size.height % 2 ? 1 : 2; + } + globalsize[0] = size.width / pxPerWorkItemX; + globalsize[1] = size.height / pxPerWorkItemY; + + // Need some padding in the private array for pixels + int privDataWidth = ROUNDUP(pxPerWorkItemX + ksize.width - 1, pxLoadNumPixels); + + // Make the global size a nice round number so the runtime can pick + // from reasonable choices for the workgroup size + const int wgRound = 256; + globalsize[0] = ROUNDUP(globalsize[0], wgRound); + + if (actual_op < 0) + actual_op = op; + + // build processing + String processing; + Mat kernel8u; + _kernel.getMat().convertTo(kernel8u, CV_8U); + for (int y = 0; y < kernel8u.rows; ++y) + for (int x = 0; x < kernel8u.cols; ++x) + if (kernel8u.at(y, x) != 0) + processing += format("PROCESS(%d,%d)", y, x); + + + static const char * const op2str[] = { "OP_ERODE", "OP_DILATE", NULL, NULL, "OP_GRADIENT", "OP_TOPHAT", "OP_BLACKHAT" }; + String opts = format("-D cn=%d " + "-D ANCHOR_X=%d -D ANCHOR_Y=%d -D KERNEL_SIZE_X=%d -D KERNEL_SIZE_Y=%d " + "-D PX_LOAD_VEC_SIZE=%d -D PX_LOAD_NUM_PX=%d -D DEPTH_%d " + "-D PX_PER_WI_X=%d -D PX_PER_WI_Y=%d -D PRIV_DATA_WIDTH=%d -D %s -D %s " + "-D PX_LOAD_X_ITERATIONS=%d -D PX_LOAD_Y_ITERATIONS=%d " + "-D srcT=%s -D srcT1=%s -D dstT=srcT -D dstT1=srcT1 -D WT=%s -D WT1=%s " + "-D convertToWT=%s -D convertToDstT=%s -D PROCESS_ELEM_=%s -D %s%s", + cn, anchor.x, anchor.y, ksize.width, ksize.height, + pxLoadVecSize, pxLoadNumPixels, depth, + pxPerWorkItemX, pxPerWorkItemY, privDataWidth, borderMap[borderType], + isolated ? "BORDER_ISOLATED" : "NO_BORDER_ISOLATED", + privDataWidth / pxLoadNumPixels, pxPerWorkItemY + ksize.height - 1, + ocl::typeToStr(type), ocl::typeToStr(depth), + haveExtraMat ? ocl::typeToStr(wtype):"srcT",//to prevent overflow - WT + haveExtraMat ? ocl::typeToStr(wdepth):"srcT1",//to prevent overflow - WT1 + haveExtraMat ? ocl::convertTypeStr(depth, wdepth, cn, cvt[0]) : "noconvert",//to prevent overflow - src to WT + haveExtraMat ? ocl::convertTypeStr(wdepth, depth, cn, cvt[1]) : "noconvert",//to prevent overflow - WT to dst + processing.c_str(), op2str[op], + actual_op == op ? "" : cv::format(" -D %s", op2str[actual_op]).c_str()); + + ocl::Kernel kernel("filterSmall", cv::ocl::imgproc::filterSmall_oclsrc, opts); + if (kernel.empty()) + return false; + + _dst.create(size, type); + UMat dst = _dst.getUMat(); + + UMat source; + if(src.u != dst.u) + source = src; + else + { + Point ofs; + int cols = src.cols, rows = src.rows; + src.locateROI(wholeSize, ofs); + src.adjustROI(ofs.y, wholeSize.height - rows - ofs.y, ofs.x, wholeSize.width - cols - ofs.x); + src.copyTo(source); + + src.adjustROI(-ofs.y, -wholeSize.height + rows + ofs.y, -ofs.x, -wholeSize.width + cols + ofs.x); + source.adjustROI(-ofs.y, -wholeSize.height + rows + ofs.y, -ofs.x, -wholeSize.width + cols + ofs.x); + source.locateROI(wholeSize, ofs); + } + + UMat extraMat = _extraMat.getUMat(); + + int idxArg = kernel.set(0, ocl::KernelArg::PtrReadOnly(source)); + idxArg = kernel.set(idxArg, (int)source.step); + int srcOffsetX = (int)((source.offset % source.step) / source.elemSize()); + int srcOffsetY = (int)(source.offset / source.step); + int srcEndX = isolated ? srcOffsetX + size.width : wholeSize.width; + int srcEndY = isolated ? srcOffsetY + size.height : wholeSize.height; + idxArg = kernel.set(idxArg, srcOffsetX); + idxArg = kernel.set(idxArg, srcOffsetY); + idxArg = kernel.set(idxArg, srcEndX); + idxArg = kernel.set(idxArg, srcEndY); + idxArg = kernel.set(idxArg, ocl::KernelArg::WriteOnly(dst)); + + if (haveExtraMat) + { + idxArg = kernel.set(idxArg, ocl::KernelArg::ReadOnlyNoSize(extraMat)); + } + + return kernel.run(2, globalsize, NULL, false); + +} + static bool ocl_morphOp(InputArray _src, OutputArray _dst, InputArray _kernel, Point anchor, int iterations, int op, int borderType, const Scalar &, int actual_op = -1, InputArray _extraMat = noArray()) { const ocl::Device & dev = ocl::Device::getDefault(); - int type = _src.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type); - bool doubleSupport = dev.doubleFPConfig() > 0; + int type = _src.type(), depth = CV_MAT_DEPTH(type), + cn = CV_MAT_CN(type), esz = CV_ELEM_SIZE(type); + Mat kernel = _kernel.getMat(); + Size ksize = kernel.data ? kernel.size() : Size(3, 3), ssize = _src.size(); + + // try to use OpenCL kernel adopted for small morph kernel + if (dev.isIntel() && !(dev.type() & ocl::Device::TYPE_CPU) && + ((ksize.width < 5 && ksize.height < 5 && esz <= 4) || + (ksize.width == 5 && ksize.height == 5 && cn == 1)) && + (iterations == 1)) + { + if (ocl_morphSmall(_src, _dst, _kernel, anchor, borderType, op, actual_op, _extraMat)) + return true; + } + bool doubleSupport = dev.doubleFPConfig() > 0; if ((depth == CV_64F && !doubleSupport) || borderType != BORDER_CONSTANT) return false; - Mat kernel = _kernel.getMat(); bool haveExtraMat = !_extraMat.empty(); - Size ksize = kernel.data ? kernel.size() : Size(3, 3), ssize = _src.size(); CV_Assert(actual_op <= 3 || haveExtraMat); if (iterations == 0 || kernel.rows*kernel.cols == 1) diff --git a/modules/imgproc/src/opencl/boxFilterSmall.cl b/modules/imgproc/src/opencl/filterSmall.cl similarity index 71% rename from modules/imgproc/src/opencl/boxFilterSmall.cl rename to modules/imgproc/src/opencl/filterSmall.cl index ff47d18..c996fb8 100755 --- a/modules/imgproc/src/opencl/boxFilterSmall.cl +++ b/modules/imgproc/src/opencl/filterSmall.cl @@ -153,35 +153,10 @@ inline bool isBorder(const struct RectCoords bounds, int2 coord, int numPixels) } #endif -inline WT getBorderPixel(const struct RectCoords bounds, int2 coord, - __global const uchar * srcptr, int srcstep) -{ -#ifdef BORDER_CONSTANT - return (WT)(0); -#else - int selected_col = coord.x; - int selected_row = coord.y; - - EXTRAPOLATE(selected_col, selected_row, - bounds.x1, bounds.y1, - bounds.x2, bounds.y2); - - __global const uchar* ptr = srcptr + mad24(selected_row, srcstep, selected_col * SRCSIZE); - return convertToWT(loadpix(ptr)); -#endif -} - -inline WT readSrcPixelSingle(int2 pos, __global const uchar * srcptr, - int srcstep, const struct RectCoords srcCoords) -{ - if (!isBorder(srcCoords, pos, 1)) - { - __global const uchar * ptr = srcptr + mad24(pos.y, srcstep, pos.x * SRCSIZE); - return convertToWT(loadpix(ptr)); - } - else - return getBorderPixel(srcCoords, pos, srcptr, srcstep); -} +#define float1 float +#define uchar1 uchar +#define int1 int +#define uint1 unit #define __CAT(x, y) x##y #define CAT(x, y) __CAT(x, y) @@ -191,7 +166,7 @@ inline WT readSrcPixelSingle(int2 pos, __global const uchar * srcptr, #define PX_LOAD_FLOAT_VEC_TYPE CAT(WT1, PX_LOAD_VEC_SIZE) #define PX_LOAD_FLOAT_VEC_CONV CAT(convert_, PX_LOAD_FLOAT_VEC_TYPE) #define PX_LOAD CAT(vload, PX_LOAD_VEC_SIZE) -#define float1 float + inline PX_LOAD_FLOAT_VEC_TYPE readSrcPixelGroup(int2 pos, __global const uchar * srcptr, int srcstep, const struct RectCoords srcCoords) @@ -218,12 +193,150 @@ inline PX_LOAD_FLOAT_VEC_TYPE readSrcPixelGroup(int2 pos, __global const uchar * #define LOOP(N, VAR, STMT) CAT(LOOP, N)((VAR), (STMT)) -__kernel void boxFilterSmall(__global const uchar * srcptr, int src_step, int srcOffsetX, int srcOffsetY, int srcEndX, int srcEndY, - __global uchar * dstptr, int dst_step, int dst_offset, int rows, int cols +#ifdef OP_BOX_FILTER +#define PROCESS_ELEM \ + WT total_sum = (WT)(0); \ + int sy = 0; \ + LOOP(KERNEL_SIZE_Y, sy, \ + { \ + int sx = 0; \ + LOOP(KERNEL_SIZE_X, sx, \ + { \ + total_sum += privateData[py + sy][px + sx]; \ + }); \ + }) + +#elif defined OP_FILTER2D + +#define DIG(a) a, +__constant WT1 kernelData[] = { COEFF }; + +#define PROCESS_ELEM \ + WT total_sum = 0; \ + int sy = 0; \ + int kernelIndex = 0; \ + LOOP(KERNEL_SIZE_Y, sy, \ + { \ + int sx = 0; \ + LOOP(KERNEL_SIZE_X, sx, \ + { \ + total_sum = fma(kernelData[kernelIndex++], privateData[py + sy][px + sx], total_sum); \ + }); \ + }) + +#elif defined OP_ERODE || defined OP_DILATE + +#ifdef DEPTH_0 +#define MIN_VAL 0 +#define MAX_VAL UCHAR_MAX +#elif defined DEPTH_1 +#define MIN_VAL SCHAR_MIN +#define MAX_VAL SCHAR_MAX +#elif defined DEPTH_2 +#define MIN_VAL 0 +#define MAX_VAL USHRT_MAX +#elif defined DEPTH_3 +#define MIN_VAL SHRT_MIN +#define MAX_VAL SHRT_MAX +#elif defined DEPTH_4 +#define MIN_VAL INT_MIN +#define MAX_VAL INT_MAX +#elif defined DEPTH_5 +#define MIN_VAL (-FLT_MAX) +#define MAX_VAL FLT_MAX +#elif defined DEPTH_6 +#define MIN_VAL (-DBL_MAX) +#define MAX_VAL DBL_MAX +#endif + +#ifdef OP_ERODE +#define VAL (WT)MAX_VAL +#elif defined OP_DILATE +#define VAL (WT)MIN_VAL +#else +#error "Unknown operation" +#endif + +#define convert_float1 convert_float +#define convert_uchar1 convert_uchar +#define convert_int1 convert_int +#define convert_uint1 convert_uint + +#ifdef OP_ERODE +#if defined INTEL_DEVICE && defined DEPTH_0 +// workaround for bug in Intel HD graphics drivers (10.18.10.3496 or older) +#define WA_CONVERT_1 CAT(convert_uint, cn) +#define WA_CONVERT_2 CAT(convert_, srcT) +#define MORPH_OP(A, B) WA_CONVERT_2(min(WA_CONVERT_1(A), WA_CONVERT_1(B))) +#else +#define MORPH_OP(A, B) min((A), (B)) +#endif +#endif +#ifdef OP_DILATE +#define MORPH_OP(A, B) max((A), (B)) +#endif + +#define PROCESS(_y, _x) \ + total_sum = convertToWT(MORPH_OP(convertToWT(total_sum), convertToWT(privateData[py + _y][px + _x]))); + +#define PROCESS_ELEM \ + WT total_sum = convertToWT(VAL); \ + PROCESS_ELEM_ + +#else +#error "No processing is specified" +#endif + +#if defined OP_GRADIENT || defined OP_TOPHAT || defined OP_BLACKHAT +#define EXTRA_PARAMS , __global const uchar * matptr, int mat_step, int mat_offset +#else +#define EXTRA_PARAMS +#endif + +inline WT getBorderPixel(const struct RectCoords bounds, int2 coord, + __global const uchar * srcptr, int srcstep) +{ +#ifdef BORDER_CONSTANT +#ifdef OP_ERODE + return (WT)(MAX_VAL); +#elif defined OP_DILATE + return (WT)(MIN_VAL); +#else + return (WT)(0); +#endif +#else + + int selected_col = coord.x; + int selected_row = coord.y; + + EXTRAPOLATE(selected_col, selected_row, + bounds.x1, bounds.y1, + bounds.x2, bounds.y2); + + __global const uchar* ptr = srcptr + mad24(selected_row, srcstep, selected_col * SRCSIZE); + return convertToWT(loadpix(ptr)); +#endif +} + +inline WT readSrcPixelSingle(int2 pos, __global const uchar * srcptr, + int srcstep, const struct RectCoords srcCoords) +{ + if (!isBorder(srcCoords, pos, 1)) + { + __global const uchar * ptr = srcptr + mad24(pos.y, srcstep, pos.x * SRCSIZE); + return convertToWT(loadpix(ptr)); + } + else + return getBorderPixel(srcCoords, pos, srcptr, srcstep); +} + + +__kernel void filterSmall(__global const uchar * srcptr, int src_step, int srcOffsetX, int srcOffsetY, int srcEndX, int srcEndY, + __global uchar * dstptr, int dst_step, int dst_offset, int rows, int cols #ifdef NORMALIZE - , float alpha + , float alpha #endif - ) + EXTRA_PARAMS ) { // for non-isolated border: offsetX, offsetY, wholeX, wholeY const struct RectCoords srcCoords = { srcOffsetX, srcOffsetY, srcEndX, srcEndY }; @@ -282,24 +395,27 @@ __kernel void boxFilterSmall(__global const uchar * srcptr, int src_step, int sr LOOP(PX_PER_WI_X, px, { int x = startX + px; - int sy = 0; - int kernelIndex = 0; - WT total_sum = (WT)(0); - - LOOP(KERNEL_SIZE_Y, sy, - { - int sx = 0; - LOOP(KERNEL_SIZE_X, sx, - { - total_sum += privateData[py + sy][px + sx]; - }); - }); - - __global dstT * dstPtr = (__global dstT *)(dstptr + mad24(y, dst_step, mad24(x, DSTSIZE, dst_offset))); + PROCESS_ELEM; + int dst_index = mad24(y, dst_step, mad24(x, DSTSIZE, dst_offset)); + __global dstT * dstPtr = (__global dstT *)(dstptr + dst_index); #ifdef NORMALIZE total_sum *= (WT)(alpha); #endif +#if defined OP_GRADIENT || defined OP_TOPHAT || defined OP_BLACKHAT + //for this type of operations SRCSIZE == DSTSIZE + int mat_index = mad24(y, mat_step, mad24(x, SRCSIZE, mat_offset)); + WT value = convertToWT(loadpix(matptr + mat_index)); + +#ifdef OP_GRADIENT + storepix(convertToDstT(convertToWT(total_sum) - convertToWT(value)), dstPtr ); +#elif defined OP_TOPHAT + storepix(convertToDstT(convertToWT(value) - convertToWT(total_sum)), dstPtr ); +#elif defined OP_BLACKHAT + storepix(convertToDstT(convertToWT(total_sum) - convertToWT(value)), dstPtr ); +#endif +#else // erode or dilate, or open-close storepix(convertToDstT(total_sum), dstPtr); +#endif }); }); } diff --git a/modules/imgproc/src/smooth.cpp b/modules/imgproc/src/smooth.cpp index 66ff429..907a659 100644 --- a/modules/imgproc/src/smooth.cpp +++ b/modules/imgproc/src/smooth.cpp @@ -720,7 +720,7 @@ static bool ocl_boxFilter( InputArray _src, OutputArray _dst, int ddepth, "-D PX_PER_WI_X=%d -D PX_PER_WI_Y=%d -D PRIV_DATA_WIDTH=%d -D %s -D %s " "-D PX_LOAD_X_ITERATIONS=%d -D PX_LOAD_Y_ITERATIONS=%d " "-D srcT=%s -D srcT1=%s -D dstT=%s -D dstT1=%s -D WT=%s -D WT1=%s " - "-D convertToWT=%s -D convertToDstT=%s%s%s", + "-D convertToWT=%s -D convertToDstT=%s%s%s -D OP_BOX_FILTER", cn, anchor.x, anchor.y, ksize.width, ksize.height, pxLoadVecSize, pxLoadNumPixels, pxPerWorkItemX, pxPerWorkItemY, privDataWidth, borderMap[borderType], @@ -734,7 +734,7 @@ static bool ocl_boxFilter( InputArray _src, OutputArray _dst, int ddepth, - if (!kernel.create("boxFilterSmall", cv::ocl::imgproc::boxFilterSmall_oclsrc, build_options)) + if (!kernel.create("filterSmall", cv::ocl::imgproc::filterSmall_oclsrc, build_options)) return false; } else diff --git a/modules/imgproc/test/ocl/test_filters.cpp b/modules/imgproc/test/ocl/test_filters.cpp index 1fe2927..aa881bd 100644 --- a/modules/imgproc/test/ocl/test_filters.cpp +++ b/modules/imgproc/test/ocl/test_filters.cpp @@ -275,14 +275,68 @@ OCL_TEST_P(Dilate, Mat) ///////////////////////////////////////////////////////////////////////////////////////////////// // MorphologyEx +IMPLEMENT_PARAM_CLASS(MorphOp, int) +PARAM_TEST_CASE(MorphologyEx, MatType, + int, // kernel size + MorphOp, // MORPH_OP + int, // iterations + bool) +{ + int type, ksize, op, iterations; + bool useRoi; + + TEST_DECLARE_INPUT_PARAMETER(src); + TEST_DECLARE_OUTPUT_PARAMETER(dst); + + virtual void SetUp() + { + type = GET_PARAM(0); + ksize = GET_PARAM(1); + op = GET_PARAM(2); + iterations = GET_PARAM(3); + useRoi = GET_PARAM(4); + } + + void random_roi(int minSize = 1) + { + if (minSize == 0) + minSize = ksize; + + Size roiSize = randomSize(minSize, MAX_VALUE); + + Border srcBorder = randomBorder(0, useRoi ? MAX_VALUE : 0); + randomSubMat(src, src_roi, roiSize, srcBorder, type, 5, 256); + + Border dstBorder = randomBorder(0, useRoi ? MAX_VALUE : 0); + randomSubMat(dst, dst_roi, roiSize, dstBorder, type, -60, 70); + + UMAT_UPLOAD_INPUT_PARAMETER(src); + UMAT_UPLOAD_OUTPUT_PARAMETER(dst); + } + + void Near() + { + int depth = CV_MAT_DEPTH(type); + bool isFP = depth >= CV_32F; -typedef FilterTestBase MorphologyEx; + if (isFP) + Near(1e-6, true); + else + Near(1, false); + } + + void Near(double threshold, bool relative) + { + if (relative) + OCL_EXPECT_MATS_NEAR_RELATIVE(dst, threshold); + else + OCL_EXPECT_MATS_NEAR(dst, threshold); + } +}; OCL_TEST_P(MorphologyEx, Mat) { Size kernelSize(ksize, ksize); - int iterations = (int)param; - int op = size.height; for (int j = 0; j < test_loop_times; j++) { @@ -377,12 +431,10 @@ OCL_INSTANTIATE_TEST_CASE_P(Filter, Dilate, Combine( OCL_INSTANTIATE_TEST_CASE_P(Filter, MorphologyEx, Combine( Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4), - Values(3, 5, 7), - Values(Size(0, 2), Size(0, 3), Size(0, 4), Size(0, 5), Size(0, 6)), // used as generator of operations - Values((BorderType)BORDER_CONSTANT), - Values(1.0, 2.0, 3.0), - Bool(), - Values(1))); // not used + Values(3, 5, 7), // kernel size + Values(MORPH_OPEN, MORPH_CLOSE, MORPH_GRADIENT, MORPH_TOPHAT, MORPH_BLACKHAT), // used as generator of operations + Values(1, 2, 3), + Bool())); } } // namespace cvtest::ocl -- 2.7.4