added cv::calcBackProject for 2-dimensional histograms
This commit is contained in:
parent
d64bea00b2
commit
b23edc34e9
@ -1942,7 +1942,7 @@ static void getUMatIndex(const std::vector<UMat> & 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<int> 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<int> 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<int> 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<int>& 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();
|
||||
|
@ -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);
|
||||
}
|
||||
}
|
||||
|
||||
|
@ -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)
|
||||
}
|
||||
};
|
||||
|
Loading…
Reference in New Issue
Block a user