From d64bea00b242ee0b5f5a87d1de476da0d603ba41 Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Mon, 23 Dec 2013 17:37:41 +0400 Subject: [PATCH] ported cv::calcBackProject to T-API --- modules/core/src/matrix.cpp | 22 +++ modules/imgproc/src/histogram.cpp | 127 ++++++++++++ .../imgproc/src/opencl/calc_back_project.cl | 133 +++++++++++++ modules/imgproc/test/ocl/test_histogram.cpp | 184 ++++++++++++++++++ 4 files changed, 466 insertions(+) create mode 100644 modules/imgproc/src/opencl/calc_back_project.cl create mode 100644 modules/imgproc/test/ocl/test_histogram.cpp diff --git a/modules/core/src/matrix.cpp b/modules/core/src/matrix.cpp index 6f2580498f..eb5d048f70 100644 --- a/modules/core/src/matrix.cpp +++ b/modules/core/src/matrix.cpp @@ -1430,6 +1430,16 @@ Size _InputArray::size(int i) const return vv[i].size(); } + if( k == STD_VECTOR_UMAT ) + { + const std::vector& vv = *(const std::vector*)obj; + if( i < 0 ) + return vv.empty() ? Size() : Size((int)vv.size(), 1); + CV_Assert( i < (int)vv.size() ); + + return vv[i].size(); + } + if( k == OPENGL_BUFFER ) { CV_Assert( i < 0 ); @@ -2262,6 +2272,12 @@ void _OutputArray::release() const return; } + if( k == UMAT ) + { + ((UMat*)obj)->release(); + return; + } + if( k == GPU_MAT ) { ((cuda::GpuMat*)obj)->release(); @@ -2301,6 +2317,12 @@ void _OutputArray::release() const return; } + if( k == STD_VECTOR_UMAT ) + { + ((std::vector*)obj)->clear(); + return; + } + CV_Error(Error::StsNotImplemented, "Unknown/unsupported array type"); } diff --git a/modules/imgproc/src/histogram.cpp b/modules/imgproc/src/histogram.cpp index 7849d5175c..2f60073bd0 100644 --- a/modules/imgproc/src/histogram.cpp +++ b/modules/imgproc/src/histogram.cpp @@ -1930,13 +1930,137 @@ void cv::calcBackProject( const Mat* images, int nimages, const int* channels, } +namespace cv { + +static void getUMatIndex(const std::vector & um, int cn, int & idx, int & cnidx) +{ + int totalChannels = 0; + for (size_t i = 0, size = um.size(); i < size; ++i) + { + int ccn = um[i].channels(); + totalChannels += ccn; + + if (totalChannels >= cn) + { + idx = i; + cnidx = i == 0 ? cn : cn % (totalChannels - ccn); + return; + } + } + + idx = cnidx = -1; +} + +static bool ocl_calcBackProject( InputArrayOfArrays _images, std::vector channels, + InputArray _hist, OutputArray _dst, + const std::vector& ranges, + float scale, size_t histdims ) +{ + const std::vector & images = *(const std::vector *)_images.getObj(); + size_t nimages = images.size(), totalcn = images[0].channels(); + + CV_Assert(nimages > 0); + Size size = images[0].size(); + int depth = images[0].depth(); + + for (size_t i = 1; i < nimages; ++i) + { + const UMat & m = images[i]; + totalcn *= m.channels(); + CV_Assert(size == m.size() && depth == m.depth()); + } + + std::sort(channels.begin(), channels.end()); + for (size_t i = 0; i < histdims; ++i) + CV_Assert(channels[i] < (int)totalcn); + + if (histdims == 1) + { + int idx, cnidx; + getUMatIndex(images, channels[0], idx, cnidx); + CV_Assert(idx >= 0); + UMat im = images[idx]; + + String opts = format("-D histdims=1 -D scn=%d", im.channels(), cnidx); + ocl::Kernel lutk("calcLUT", ocl::imgproc::calc_back_project_oclsrc, opts); + if (lutk.empty()) + return false; + + size_t lsize = 256; + UMat lut(1, (int)lsize, CV_32SC1), hist = _hist.getUMat(), uranges(ranges, true); + + lutk.args(ocl::KernelArg::ReadOnlyNoSize(hist), hist.rows, + ocl::KernelArg::PtrWriteOnly(lut), scale, ocl::KernelArg::PtrReadOnly(uranges)); + if (!lutk.run(1, &lsize, NULL, false)) + return false; + + ocl::Kernel mapk("LUT", ocl::imgproc::calc_back_project_oclsrc, opts); + if (mapk.empty()) + return false; + + _dst.create(size, depth); + UMat dst = _dst.getUMat(); + + im.offset += cnidx; + mapk.args(ocl::KernelArg::ReadOnlyNoSize(im), ocl::KernelArg::PtrReadOnly(lut), + ocl::KernelArg::WriteOnly(dst)); + + size_t globalsize[2] = { size.width, size.height }; + return mapk.run(2, globalsize, NULL, false); + } + else if (histdims == 2) + { + int idx0, idx1, cnidx0, cnidx1; + getUMatIndex(images, channels[0], idx0, cnidx0); + getUMatIndex(images, channels[1], idx1, cnidx1); + printf("%d) channels = %d, indx = %d, cnidx = %d\n", images[0].channels(), channels[0], idx0, cnidx0); + printf("%d) channels = %d, indx = %d, cnidx = %d\n", images[1].channels(), channels[1], idx1, cnidx1); + CV_Assert(idx0 >= 0 && idx1 >= 0); + UMat im0 = images[idx0], im1 = images[idx1]; + + String opts = format("-D histdims=2 -D scn0=%d -D scn1=%d", + im0.channels(), im1.channels()); + ocl::Kernel k("calcBackProject", ocl::imgproc::calc_back_project_oclsrc, opts); + if (k.empty()) + return false; + + _dst.create(size, depth); + UMat dst = _dst.getUMat(), hist = _hist.getUMat(), uranges(ranges, true); + + im0.offset += cnidx0; + im1.offset += cnidx1; + k.args(ocl::KernelArg::ReadOnlyNoSize(im0), ocl::KernelArg::ReadOnlyNoSize(im1), + ocl::KernelArg::ReadOnly(hist), ocl::KernelArg::WriteOnly(dst), scale, + ocl::KernelArg::PtrReadOnly(uranges)); + + size_t globalsize[2] = { size.width, size.height }; + return k.run(2, globalsize, NULL, false); + } + return false; +} + +} + void cv::calcBackProject( InputArrayOfArrays images, const std::vector& channels, InputArray hist, OutputArray dst, const std::vector& ranges, double scale ) { + Size histSize = hist.size(); + bool _1D = histSize.height == 1 || histSize.width == 1; + size_t histdims = _1D ? 1 : hist.dims(); + + if (ocl::useOpenCL() && images.isUMatVector() && dst.isUMat() && hist.type() == CV_32FC1 && + histdims <= 2 && ranges.size() == histdims * 2 && histdims == channels.size() /*&& + ocl_calcBackProject(images, channels, hist, dst, ranges, scale)*/) + { + CV_Assert(ocl_calcBackProject(images, channels, hist, dst, ranges, (float)scale, histdims)); + return; + } + Mat H0 = hist.getMat(), H; int hcn = H0.channels(); + if( hcn > 1 ) { CV_Assert( H0.isContinuous() ); @@ -1947,12 +2071,15 @@ void cv::calcBackProject( InputArrayOfArrays images, const std::vector& cha } else H = H0; + bool _1d = H.rows == 1 || H.cols == 1; int i, dims = H.dims, rsz = (int)ranges.size(), csz = (int)channels.size(); int nimages = (int)images.total(); + CV_Assert(nimages > 0); CV_Assert(rsz == dims*2 || (rsz == 2 && _1d) || (rsz == 0 && images.depth(0) == CV_8U)); CV_Assert(csz == 0 || csz == dims || (csz == 1 && _1d)); + float* _ranges[CV_MAX_DIM]; if( rsz > 0 ) { diff --git a/modules/imgproc/src/opencl/calc_back_project.cl b/modules/imgproc/src/opencl/calc_back_project.cl new file mode 100644 index 0000000000..b5b0c03a25 --- /dev/null +++ b/modules/imgproc/src/opencl/calc_back_project.cl @@ -0,0 +1,133 @@ +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved. +// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. +// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// @Authors +// Niko Li, newlife20080214@gmail.com +// Jia Haipeng, jiahaipeng95@gmail.com +// Xu Pang, pangxu010@163.com +// Wenju He, wenju@multicorewareinc.com +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors as is and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +// + +#if histdims == 1 + +#define OUT_OF_RANGE -1 + +__kernel void calcLUT(__global const uchar * histptr, int hist_step, int hist_offset, int hist_bins, + __global int * lut, float scale, __constant float * ranges) +{ + int x = get_global_id(0); + float value = convert_float(x); + + if (value > ranges[1] || value < ranges[0]) + lut[x] = OUT_OF_RANGE; + else + { + float lb = ranges[0], ub = ranges[1], gap = (ub - lb) / hist_bins; + value -= lb; + int bin = convert_int_sat_rtn(value / gap); + + if (bin >= hist_bins) + lut[x] = OUT_OF_RANGE; + else + { + int hist_index = mad24(hist_step, bin, hist_offset); + __global const float * hist = (__global const float *)(histptr + hist_index); + + lut[x] = (int)convert_uchar_sat_rte(hist[0] * scale); + } + } +} + +__kernel void LUT(__global const uchar * src, int src_step, int src_offset, + __global const int * lut, + __global uchar * dst, 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 src_index = mad24(y, src_step, src_offset + x * scn); + int dst_index = mad24(y, dst_step, dst_offset + x); + + int value = lut[src[src_index]]; + dst[dst_index] = value == OUT_OF_RANGE ? 0 : convert_uchar(value); + } +} + +#elif histdims == 2 + +#define OUT_OF_RANGES(i) ( (value##i > ranges[(i<<1)+1]) || (value##i < ranges[i<<1]) ) +#define CALCULATE_BIN(i) \ + float lb##i = ranges[i<<1], ub##i = ranges[(i<<1)+1], gap##i = (ub##i - lb##i) / hist_bins##i; \ + value##i -= ranges[i<<1]; \ + int bin##i = convert_int_sat_rtn(value##i / gap##i) + +__kernel void calcBackProject(__global const uchar * src0, int src0_step, int src0_offset, + __global const uchar * src1, int src1_step, int src1_offset, + __global const uchar * histptr, int hist_step, int hist_offset, int hist_bins0, int hist_bins1, + __global uchar * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols, + float scale, __constant float * ranges) +{ + int x = get_global_id(0); + int y = get_global_id(1); + + if (x < dst_cols && y < dst_rows) + { + int src0_index = mad24(src0_step, y, src0_offset + x * scn0); + int src1_index = mad24(src1_step, y, src1_offset + x * scn1); + int dst_index = mad24(dst_step, y, dst_offset + x); + + float value0 = convert_float(src0[src0_index]), value1 = convert_float(src1[src1_index]); + if (OUT_OF_RANGES(0) || OUT_OF_RANGES(1)) + dst[dst_index] = 0; + else + { + CALCULATE_BIN(0); + CALCULATE_BIN(1); + + if (bin0 >= hist_bins0 || bin1 >= hist_bins1) + dst[dst_index] = 0; + else + { + int hist_index = mad24(hist_step, bin0, hist_offset + bin1 * (int)sizeof(float)); + __global const float * hist = (__global const float *)(histptr + hist_index); + + dst[dst_index] = convert_uchar_sat_rte(scale * hist[0]); + } + } + } +} + +#else +#error "(nimages <= 2) should be true" +#endif diff --git a/modules/imgproc/test/ocl/test_histogram.cpp b/modules/imgproc/test/ocl/test_histogram.cpp new file mode 100644 index 0000000000..6714909ace --- /dev/null +++ b/modules/imgproc/test/ocl/test_histogram.cpp @@ -0,0 +1,184 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved. +// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. +// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// @Authors +// Niko Li, newlife20080214@gmail.com +// Jia Haipeng, jiahaipeng95@gmail.com +// Shengen Yan, yanshengen@gmail.com +// Jiang Liyuan, lyuan001.good@163.com +// Rock Li, Rock.Li@amd.com +// Wu Zailong, bullet@yeah.net +// Xu Pang, pangxu010@163.com +// Sen Liu, swjtuls1987@126.com +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors "as is" and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ + +#include "test_precomp.hpp" +#include "cvconfig.h" +#include "opencv2/ts/ocl_test.hpp" + +#ifdef HAVE_OPENCL + +namespace cvtest { +namespace ocl { + +/////////////////////////////////////////////////////////////////////////////// + +PARAM_TEST_CASE(CalcBackProject, MatDepth, int, bool) +{ + int depth, N; + bool useRoi; + + std::vector ranges; + std::vector channels; + double scale; + + std::vector images; + std::vector images_roi; + std::vector uimages; + std::vector uimages_roi; + + TEST_DECLARE_INPUT_PARAMETER(hist) + TEST_DECLARE_OUTPUT_PARAMETER(dst) + + virtual void SetUp() + { + depth = GET_PARAM(0); + N = GET_PARAM(1); + useRoi = GET_PARAM(2); + + ASSERT_GE(2, N); + + images.resize(N); + images_roi.resize(N); + uimages.resize(N); + uimages_roi.resize(N); + } + + virtual void random_roi() + { + Size roiSize = randomSize(1, MAX_VALUE); + + int totalChannels = 0; + for (int i = 0; i < N; ++i) + { + Border srcBorder = randomBorder(0, useRoi ? MAX_VALUE : 0); + int cn = randomInt(1, 5); + randomSubMat(images[i], images_roi[i], roiSize, srcBorder, CV_MAKE_TYPE(depth, cn), 0, 125); + + ranges.push_back(10); + ranges.push_back(100); + + channels.push_back(randomInt(0, cn) + totalChannels); + totalChannels += cn; + } + + Mat tmpHist; + { + std::vector hist_size(N); + for (int i = 0 ; i < N; ++i) + hist_size[i] = randomInt(10, 50); + + cv::calcHist(images_roi, channels, noArray(), tmpHist, hist_size, ranges); + ASSERT_EQ(CV_32FC1, tmpHist.type()); + } + + Border histBorder = randomBorder(0, useRoi ? MAX_VALUE : 0); + randomSubMat(hist, hist_roi, tmpHist.size(), histBorder, tmpHist.type(), 0, MAX_VALUE); + tmpHist.copyTo(hist_roi); + + Border dstBorder = randomBorder(0, useRoi ? MAX_VALUE : 0); + randomSubMat(dst, dst_roi, roiSize, dstBorder, CV_MAKE_TYPE(depth, 1), 5, 16); + + for (int i = 0; i < N; ++i) + { + images[i].copyTo(uimages[i]); + + Size _wholeSize; + Point ofs; + images_roi[i].locateROI(_wholeSize, ofs); + + uimages_roi[i] = uimages[i](Rect(ofs.x, ofs.y, images_roi[i].cols, images_roi[i].rows)); + } + + UMAT_UPLOAD_INPUT_PARAMETER(hist) + UMAT_UPLOAD_OUTPUT_PARAMETER(dst) + + scale = randomDouble(0.1, 1); + } + + void Near() + { +// std::cout << "Src: " << std::endl << src_roi[0] << std::endl; +// std::cout << "Hist: " << std::endl << hist_roi << std::endl; + std::cout << "OpenCV: " << std::endl << dst_roi << std::endl; + std::cout << "OpenCL: " << std::endl << udst_roi.getMat(ACCESS_READ) << std::endl; + + Mat diff; + cv::absdiff(dst_roi, udst_roi, diff); + std::cout << "Difference: " << std::endl << diff << std::endl; + + OCL_EXPECT_MATS_NEAR(dst, 0.0) + } +}; + +//////////////////////////////// CalcBackProject ////////////////////////////////////////////// + +OCL_TEST_P(CalcBackProject, Mat) +{ + for (int j = 0; j < test_loop_times; j++) + { + random_roi(); + + OCL_OFF(cv::calcBackProject(images_roi, channels, hist_roi, dst_roi, ranges, scale)); + OCL_ON(cv::calcBackProject(uimages_roi, channels, uhist_roi, udst_roi, ranges, scale)); + + Near(); + } +} + +///////////////////////////////////////////////////////////////////////////////////// + +OCL_INSTANTIATE_TEST_CASE_P(Imgproc, CalcBackProject, Combine(Values((MatDepth)CV_8U), Values(1, 2), Bool())); + +} } // namespace cvtest::ocl + +#endif // HAVE_OPENCL