diff --git a/modules/imgproc/src/histogram.cpp b/modules/imgproc/src/histogram.cpp index 2f60073bd..1aee957b8 100644 --- a/modules/imgproc/src/histogram.cpp +++ b/modules/imgproc/src/histogram.cpp @@ -1942,7 +1942,7 @@ static void getUMatIndex(const std::vector & um, int cn, int & idx, int & if (totalChannels >= cn) { - idx = i; + idx = (int)i; cnidx = i == 0 ? cn : cn % (totalChannels - ccn); return; } @@ -1966,7 +1966,7 @@ static bool ocl_calcBackProject( InputArrayOfArrays _images, std::vector ch for (size_t i = 1; i < nimages; ++i) { const UMat & m = images[i]; - totalcn *= m.channels(); + totalcn += m.channels(); CV_Assert(size == m.size() && depth == m.depth()); } @@ -1981,7 +1981,7 @@ static bool ocl_calcBackProject( InputArrayOfArrays _images, std::vector ch CV_Assert(idx >= 0); UMat im = images[idx]; - String opts = format("-D histdims=1 -D scn=%d", im.channels(), cnidx); + String opts = format("-D histdims=1 -D scn=%d", im.channels()); ocl::Kernel lutk("calcLUT", ocl::imgproc::calc_back_project_oclsrc, opts); if (lutk.empty()) return false; @@ -2013,28 +2013,47 @@ static bool ocl_calcBackProject( InputArrayOfArrays _images, std::vector ch 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()) + // Lut for the first dimension + String opts = format("-D histdims=2 -D scn1=%d -D scn2=%d", im0.channels(), im1.channels()); + ocl::Kernel lutk1("calcLUT", ocl::imgproc::calc_back_project_oclsrc, opts); + if (lutk1.empty()) + return false; + + size_t lsize = 256; + UMat lut(1, (int)lsize<<1, CV_32SC1), uranges(ranges, true), hist = _hist.getUMat(); + + lutk1.args(hist.rows, ocl::KernelArg::PtrWriteOnly(lut), (int)0, ocl::KernelArg::PtrReadOnly(uranges), (int)0); + if (!lutk1.run(1, &lsize, NULL, false)) + return false; + + // lut for the second dimension + ocl::Kernel lutk2("calcLUT", ocl::imgproc::calc_back_project_oclsrc, opts); + if (lutk2.empty()) + return false; + + lut.offset += lsize * sizeof(int); + lutk2.args(hist.cols, ocl::KernelArg::PtrWriteOnly(lut), (int)256, ocl::KernelArg::PtrReadOnly(uranges), (int)2); + if (!lutk2.run(1, &lsize, NULL, false)) + return false; + + // perform lut + ocl::Kernel mapk("LUT", ocl::imgproc::calc_back_project_oclsrc, opts); + if (mapk.empty()) return false; _dst.create(size, depth); - UMat dst = _dst.getUMat(), hist = _hist.getUMat(), uranges(ranges, true); + UMat dst = _dst.getUMat(); 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)); + mapk.args(ocl::KernelArg::ReadOnlyNoSize(im0), ocl::KernelArg::ReadOnlyNoSize(im1), + ocl::KernelArg::ReadOnlyNoSize(hist), ocl::KernelArg::PtrReadOnly(lut), scale, ocl::KernelArg::WriteOnly(dst)); size_t globalsize[2] = { size.width, size.height }; - return k.run(2, globalsize, NULL, false); + return mapk.run(2, globalsize, NULL, false); } return false; } @@ -2051,12 +2070,9 @@ void cv::calcBackProject( InputArrayOfArrays images, const std::vector& cha 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)); + histdims <= 2 && ranges.size() == histdims * 2 && histdims == channels.size() && + ocl_calcBackProject(images, channels, hist, dst, ranges, (float)scale, histdims)) return; - } Mat H0 = hist.getMat(), H; int hcn = H0.channels(); diff --git a/modules/imgproc/src/opencl/calc_back_project.cl b/modules/imgproc/src/opencl/calc_back_project.cl index b5b0c03a2..ec9247154 100644 --- a/modules/imgproc/src/opencl/calc_back_project.cl +++ b/modules/imgproc/src/opencl/calc_back_project.cl @@ -37,10 +37,10 @@ // // -#if histdims == 1 - #define OUT_OF_RANGE -1 +#if histdims == 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) { @@ -68,7 +68,7 @@ __kernel void calcLUT(__global const uchar * histptr, int hist_step, int hist_of } __kernel void LUT(__global const uchar * src, int src_step, int src_offset, - __global const int * lut, + __constant int * lut, __global uchar * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols) { int x = get_global_id(0); @@ -86,45 +86,47 @@ __kernel void LUT(__global const uchar * src, int src_step, int src_offset, #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 calcLUT(int hist_bins, __global int * lut, int lut_offset, + __constant float * ranges, int roffset) +{ + int x = get_global_id(0); + float value = convert_float(x); -__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) + ranges += roffset; + lut += lut_offset; + + 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); + + lut[x] = bin >= hist_bins ? OUT_OF_RANGE : bin; + } +} + +__kernel void LUT(__global const uchar * src1, int src1_step, int src1_offset, + __global const uchar * src2, int src2_step, int src2_offset, + __global const uchar * histptr, int hist_step, int hist_offset, + __constant int * lut, float scale, + __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 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); + int src1_index = mad24(y, src1_step, src1_offset + x * scn1); + int src2_index = mad24(y, src2_step, src2_offset + x * scn2); + int dst_index = mad24(y, dst_step, 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]); - } - } + int bin1 = lut[src1[src1_index]]; + int bin2 = lut[src2[src2_index] + 256]; + dst[dst_index] = bin1 == OUT_OF_RANGE || bin2 == OUT_OF_RANGE ? 0 : + convert_uchar_sat_rte(*(__global const float *)(histptr + + mad24(hist_step, bin1, hist_offset + bin2 * (int)sizeof(float))) * scale); } } diff --git a/modules/imgproc/test/ocl/test_histogram.cpp b/modules/imgproc/test/ocl/test_histogram.cpp index 6714909ac..d6cf6efa1 100644 --- a/modules/imgproc/test/ocl/test_histogram.cpp +++ b/modules/imgproc/test/ocl/test_histogram.cpp @@ -147,15 +147,6 @@ PARAM_TEST_CASE(CalcBackProject, MatDepth, int, bool) 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) } };