ported cv::calcBackProject to T-API

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<UMat>& vv = *(const std::vector<UMat>*)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<UMat>*)obj)->clear();
+        return;
+    }
+
     CV_Error(Error::StsNotImplemented, "Unknown/unsupported array type");
 }
 

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<UMat> & 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<int> channels,
+                                 InputArray _hist, OutputArray _dst,
+                                 const std::vector<float>& ranges,
+                                 float scale, size_t histdims )
+{
+    const std::vector<UMat> & images = *(const std::vector<UMat> *)_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<int>& channels,
                           InputArray hist, OutputArray dst,
                           const std::vector<float>& 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<int>& 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 )
     {

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

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<float> ranges;
+    std::vector<int> channels;
+    double scale;
+
+    std::vector<Mat> images;
+    std::vector<Mat> images_roi;
+    std::vector<UMat> uimages;
+    std::vector<UMat> 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<int> 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