#define DATA_SIZE ((int)sizeof(type))
#define ELEM_TYPE elem_type
#define ELEM_SIZE ((int)sizeof(elem_type))
-#define CN cn
-#define SQSUMS_PTR(ox, oy) mad24(gidy + oy, img_sqsums_step, gidx*CN + img_sqsums_offset + ox*CN)
-#define SUMS_PTR(ox, oy) mad24(gidy + oy, img_sums_step, gidx*CN + img_sums_offset + ox*CN)
+#define SQSUMS_PTR(ox, oy) mad24(y + oy, src_sqsums_step, mad24(x + ox, cn, src_sqsums_offset))
+#define SUMS_PTR(ox, oy) mad24(y + oy, src_sums_step, mad24(x + ox, cn, src_sums_offset))
inline float normAcc(float num, float denum)
{
- if(fabs(num) < denum)
- {
+ if (fabs(num) < denum)
return num / denum;
- }
- if(fabs(num) < denum * 1.125f)
- {
+ if (fabs(num) < denum * 1.125f)
return num > 0 ? 1 : -1;
- }
return 0;
}
inline float normAcc_SQDIFF(float num, float denum)
{
- if(fabs(num) < denum)
- {
+ if (fabs(num) < denum)
return num / denum;
- }
- if(fabs(num) < denum * 1.125f)
- {
+ if (fabs(num) < denum * 1.125f)
return num > 0 ? 1 : -1;
- }
return 1;
}
-//////////////////////////////////////////CCORR/////////////////////////////////////////////////////////////////////////
+#define noconvert
-__kernel void matchTemplate_Naive_CCORR (__global const uchar * img,int img_step,int img_offset,
- __global const uchar * tpl,int tpl_step,int tpl_offset,int tpl_rows, int tpl_cols,
- __global uchar * res,int res_step,int res_offset,int res_rows,int res_cols)
+#if cn == 1
+#define convertToDT(value) (float)(value)
+#elif cn == 2
+#define convertToDT(value) (float)(value.x + value.y)
+#elif cn == 4
+#define convertToDT(value) (float)(value.x + value.y + value.z + value.w)
+#else
+#error "cn should be 1, 2 or 4"
+#endif
+
+#ifdef CALC_SUM
+
+__kernel void calcSum(__global const uchar * templateptr, int template_step, int template_offset,
+ int template_rows, int template_cols, __global float * result)
{
- int gidx = get_global_id(0);
- int gidy = get_global_id(1);
- int i,j;
- float sum = 0;
+ __global const T * template = (__global const T *)(templateptr + template_offset);
- int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float));
+ WT res = (WT)(0);
- if(gidx < res_cols && gidy < res_rows)
+ for (int y = 0; y < template_rows; ++y)
{
- for(i = 0; i < tpl_rows; i ++)
+ for (int x = 0; x < template_cols; ++x)
{
- __global const ELEM_TYPE * img_ptr = (__global const ELEM_TYPE *)(img + mad24(gidy + i, img_step, gidx*DATA_SIZE + img_offset));
- __global const ELEM_TYPE * tpl_ptr = (__global const ELEM_TYPE *)(tpl + mad24(i, tpl_step, tpl_offset));
+ WT value = convertToWT(template[x]);
+#ifdef SUM_2
+#if wdepth == 4
+ res = mad24(value, value, res);
+#else
+ res = mad(value, value, res);
+#endif
+#elif defined SUM_1
+ res += value;
+#else
+#error "No operation is specified"
+#endif
+ }
+
+ template = (__global const T *)((__global const uchar *)template + template_step);
+ }
- for(j = 0; j < tpl_cols; j ++)
+ result[0] = convertToDT(res);
+}
-#pragma unroll
- for (int c = 0; c < CN; c++)
+#elif defined CCORR
- sum += (float)(img_ptr[j*CN+c] * tpl_ptr[j*CN+c]);
+__kernel void matchTemplate_Naive_CCORR(__global const uchar * srcptr, int src_step, int src_offset,
+ __global const uchar * templateptr, int template_step, int template_offset, int template_rows, int template_cols,
+ __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)
+ {
+ WT sum = (WT)(0);
+
+ __global const T * src = (__global const T *)(srcptr + mad24(y, src_step, mad24(x, (int)sizeof(T), src_offset)));
+ __global const T * template = (__global const T *)(templateptr + template_offset);
+
+ for (int i = 0; i < template_rows; ++i)
+ {
+ for (int j = 0; j < template_cols; ++j)
+#if wdepth == 4
+ sum = mad24(convertToWT(src[j]), convertToWT(template[j]), sum);
+#else
+ sum = mad(convertToWT(src[j]), convertToWT(template[j]), sum);
+#endif
+
+ src = (__global const T *)((__global const uchar *)src + src_step);
+ template = (__global const T *)((__global const uchar *)template + template_step);
}
- __global float * result = (__global float *)(res+res_idx);
- *result = sum;
+
+ int dst_idx = mad24(y, dst_step, mad24(x, (int)sizeof(float), dst_offset));
+ *(__global float *)(dst + dst_idx) = convertToDT(sum);
}
}
-__kernel void matchTemplate_CCORR_NORMED ( __global const uchar * img_sqsums, int img_sqsums_step, int img_sqsums_offset,
- __global uchar * res, int res_step, int res_offset, int res_rows, int res_cols,
- int tpl_rows, int tpl_cols, float tpl_sqsum)
-{
- int gidx = get_global_id(0);
- int gidy = get_global_id(1);
-
- img_sqsums_step /= sizeof(float);
- img_sqsums_offset /= sizeof(float);
+#elif defined CCORR_NORMED
- int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float));
+__kernel void matchTemplate_CCORR_NORMED(__global const uchar * src_sqsums, int src_sqsums_step, int src_sqsums_offset,
+ __global uchar * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols,
+ int template_rows, int template_cols, __global const float * template_sqsum)
+{
+ int x = get_global_id(0);
+ int y = get_global_id(1);
- if(gidx < res_cols && gidy < res_rows)
+ if (x < dst_cols && y < dst_rows)
{
- __global float * sqsum = (__global float*)(img_sqsums);
- float image_sqsum_ = (float)(
- (sqsum[SQSUMS_PTR(tpl_cols, tpl_rows)] - sqsum[SQSUMS_PTR(tpl_cols, 0)]) -
- (sqsum[SQSUMS_PTR(0, tpl_rows)] - sqsum[SQSUMS_PTR(0, 0)]));
+ __global const float * sqsum = (__global const float *)(src_sqsums);
- __global float * result = (__global float *)(res+res_idx);
- *result = normAcc(*result, sqrt(image_sqsum_ * tpl_sqsum));
+ src_sqsums_step /= sizeof(float);
+ src_sqsums_offset /= sizeof(float);
+ float image_sqsum_ = (float)(sqsum[SQSUMS_PTR(template_cols, template_rows)] - sqsum[SQSUMS_PTR(template_cols, 0)] -
+ sqsum[SQSUMS_PTR(0, template_rows)] + sqsum[SQSUMS_PTR(0, 0)]);
+
+ int dst_idx = mad24(y, dst_step, mad24(x, (int)sizeof(float), dst_offset));
+ __global float * dstult = (__global float *)(dst + dst_idx);
+ *dstult = normAcc(*dstult, sqrt(image_sqsum_ * template_sqsum[0]));
}
}
-////////////////////////////////////////////SQDIFF////////////////////////////////////////////////////////////////////////
+#elif defined SQDIFF
-__kernel void matchTemplate_Naive_SQDIFF(__global const uchar * img,int img_step,int img_offset,
- __global const uchar * tpl,int tpl_step,int tpl_offset,int tpl_rows, int tpl_cols,
- __global uchar * res,int res_step,int res_offset,int res_rows,int res_cols)
+__kernel void matchTemplate_Naive_SQDIFF(__global const uchar * srcptr, int src_step, int src_offset,
+ __global const uchar * templateptr, int template_step, int template_offset, int template_rows, int template_cols,
+ __global uchar * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols)
{
- int gidx = get_global_id(0);
- int gidy = get_global_id(1);
- int i,j;
- float delta;
- float sum = 0;
-
- int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float));
+ int x = get_global_id(0);
+ int y = get_global_id(1);
- if(gidx < res_cols && gidy < res_rows)
+ if (x < dst_cols && y < dst_rows)
{
- for(i = 0; i < tpl_rows; i ++)
- {
- __global const ELEM_TYPE * img_ptr = (__global const ELEM_TYPE *)(img + mad24(gidy + i, img_step, gidx*DATA_SIZE + img_offset));
- __global const ELEM_TYPE * tpl_ptr = (__global const ELEM_TYPE *)(tpl + mad24(i, tpl_step, tpl_offset));
+ __global const T * src = (__global const T *)(srcptr + mad24(y, src_step, mad24(x, (int)sizeof(T), src_offset)));
+ __global const T * template = (__global const T *)(templateptr + template_offset);
- for(j = 0; j < tpl_cols; j ++)
+ WT sum = (WT)(0), value;
-#pragma unroll
- for (int c = 0; c < CN; c++)
- {
- delta = (float)(img_ptr[j*CN+c] - tpl_ptr[j*CN+c]);
- sum += delta*delta;
- }
+ for (int i = 0; i < template_rows; ++i)
+ {
+ for (int j = 0; j < template_cols; ++j)
+ {
+ value = convertToWT(src[j]) - convertToWT(template[j]);
+#if wdepth == 4
+ sum = mad24(value, value, sum);
+#else
+ sum = mad(value, value, sum);
+#endif
+ }
+
+ src = (__global const T *)((__global const uchar *)src + src_step);
+ template = (__global const T *)((__global const uchar *)template + template_step);
}
- __global float * result = (__global float *)(res+res_idx);
- *result = sum;
+
+ int dst_idx = mad24(y, dst_step, mad24(x, (int)sizeof(float), dst_offset));
+ *(__global float *)(dst + dst_idx) = convertToDT(sum);
}
}
-__kernel void matchTemplate_SQDIFF_NORMED ( __global const uchar * img_sqsums, int img_sqsums_step, int img_sqsums_offset,
- __global uchar * res, int res_step, int res_offset, int res_rows, int res_cols,
- int tpl_rows, int tpl_cols, float tpl_sqsum)
-{
- int gidx = get_global_id(0);
- int gidy = get_global_id(1);
-
- img_sqsums_step /= sizeof(float);
- img_sqsums_offset /= sizeof(float);
+#elif defined SQDIFF_NORMED
- int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float));
+__kernel void matchTemplate_SQDIFF_NORMED(__global const uchar * src_sqsums, int src_sqsums_step, int src_sqsums_offset,
+ __global uchar * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols,
+ int template_rows, int template_cols, __global const float * template_sqsum)
+{
+ int x = get_global_id(0);
+ int y = get_global_id(1);
- if(gidx < res_cols && gidy < res_rows)
+ if (x < dst_cols && y < dst_rows)
{
- __global float * sqsum = (__global float*)(img_sqsums);
- float image_sqsum_ = (float)(
- (sqsum[SQSUMS_PTR(tpl_cols, tpl_rows)] - sqsum[SQSUMS_PTR(tpl_cols, 0)]) -
- (sqsum[SQSUMS_PTR(0, tpl_rows)] - sqsum[SQSUMS_PTR(0, 0)]));
+ src_sqsums_step /= sizeof(float);
+ src_sqsums_offset /= sizeof(float);
- __global float * result = (__global float *)(res+res_idx);
+ __global const float * sqsum = (__global const float *)(src_sqsums);
+ float image_sqsum_ = (float)(
+ (sqsum[SQSUMS_PTR(template_cols, template_rows)] - sqsum[SQSUMS_PTR(template_cols, 0)]) -
+ (sqsum[SQSUMS_PTR(0, template_rows)] - sqsum[SQSUMS_PTR(0, 0)]));
+ float template_sqsum_value = template_sqsum[0];
- *result = normAcc_SQDIFF(image_sqsum_ - 2.f * result[0] + tpl_sqsum, sqrt(image_sqsum_ * tpl_sqsum));
+ int dst_idx = mad24(y, dst_step, mad24(x, (int)sizeof(float), dst_offset));
+ __global float * dstult = (__global float *)(dst + dst_idx);
+ *dstult = normAcc_SQDIFF(image_sqsum_ - 2.0f * dstult[0] + template_sqsum_value, sqrt(image_sqsum_ * template_sqsum_value));
}
}
-////////////////////////////////////////////CCOEFF/////////////////////////////////////////////////////////////////
-
-__kernel void matchTemplate_Prepared_CCOEFF_C1 (__global const uchar * img_sums, int img_sums_step, int img_sums_offset,
- __global uchar * res, int res_step, int res_offset, int res_rows, int res_cols,
- int tpl_rows, int tpl_cols, float tpl_sum)
-{
- int gidx = get_global_id(0);
- int gidy = get_global_id(1);
+#elif defined CCOEFF
- img_sums_step /= ELEM_SIZE;
- img_sums_offset /= ELEM_SIZE;
+#if cn == 1
- int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float));
- float image_sum_ = 0;
+__kernel void matchTemplate_Prepared_CCOEFF(__global const uchar * src_sums, int src_sums_step, int src_sums_offset,
+ __global uchar * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols,
+ int template_rows, int template_cols, float template_sum)
+{
+ int x = get_global_id(0);
+ int y = get_global_id(1);
- if(gidx < res_cols && gidy < res_rows)
+ if (x < dst_cols && y < dst_rows)
{
- __global ELEM_TYPE* sum = (__global ELEM_TYPE*)(img_sums);
+ __global ELEM_TYPE* sum = (__global ELEM_TYPE*)(src_sums);
- image_sum_ += (float)((sum[SUMS_PTR(tpl_cols, tpl_rows)] - sum[SUMS_PTR(tpl_cols, 0)])-
- (sum[SUMS_PTR(0, tpl_rows)] - sum[SUMS_PTR(0, 0)])) * tpl_sum;
+ src_sums_step /= ELEM_SIZE;
+ src_sums_offset /= ELEM_SIZE;
+ float image_sum_ = (float)((sum[SUMS_PTR(template_cols, template_rows)] - sum[SUMS_PTR(template_cols, 0)])-
+ (sum[SUMS_PTR(0, template_rows)] - sum[SUMS_PTR(0, 0)])) * template_sum;
- __global float * result = (__global float *)(res+res_idx);
- *result -= image_sum_;
+ int dst_idx = mad24(y, dst_step, mad24(x, (int)sizeof(float), dst_offset));
+ __global float * dstult = (__global float *)(dst + dst_idx);
+ *dstult -= image_sum_;
}
}
-__kernel void matchTemplate_Prepared_CCOEFF_C2 (__global const uchar * img_sums, int img_sums_step, int img_sums_offset,
- __global uchar * res, int res_step, int res_offset, int res_rows, int res_cols,
- int tpl_rows, int tpl_cols, float tpl_sum_0,float tpl_sum_1)
-{
- int gidx = get_global_id(0);
- int gidy = get_global_id(1);
-
- img_sums_step /= ELEM_SIZE;
- img_sums_offset /= ELEM_SIZE;
+#elif cn == 2
- int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float));
- float image_sum_ = 0;
+__kernel void matchTemplate_Prepared_CCOEFF(__global const uchar * src_sums, int src_sums_step, int src_sums_offset,
+ __global uchar * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols,
+ int template_rows, int template_cols, float template_sum_0, float template_sum_1)
+{
+ int x = get_global_id(0);
+ int y = get_global_id(1);
- if(gidx < res_cols && gidy < res_rows)
+ if (x < dst_cols && y < dst_rows)
{
- __global ELEM_TYPE* sum = (__global ELEM_TYPE*)(img_sums);
+ src_sums_step /= ELEM_SIZE;
+ src_sums_offset /= ELEM_SIZE;
+
+ __global ELEM_TYPE* sum = (__global ELEM_TYPE*)(src_sums);
- image_sum_ += tpl_sum_0 * (float)((sum[SUMS_PTR(tpl_cols, tpl_rows)] - sum[SUMS_PTR(tpl_cols, 0)]) -(sum[SUMS_PTR(0, tpl_rows)] - sum[SUMS_PTR(0, 0)]));
- image_sum_ += tpl_sum_1 * (float)((sum[SUMS_PTR(tpl_cols, tpl_rows)+1] - sum[SUMS_PTR(tpl_cols, 0)+1])-(sum[SUMS_PTR(0, tpl_rows)+1] - sum[SUMS_PTR(0, 0)+1]));
+ float image_sum_ = template_sum_0 * (float)((sum[SUMS_PTR(template_cols, template_rows)] - sum[SUMS_PTR(template_cols, 0)]) -(sum[SUMS_PTR(0, template_rows)] - sum[SUMS_PTR(0, 0)]));
+ image_sum_ += template_sum_1 * (float)((sum[SUMS_PTR(template_cols, template_rows)+1] - sum[SUMS_PTR(template_cols, 0)+1])-(sum[SUMS_PTR(0, template_rows)+1] - sum[SUMS_PTR(0, 0)+1]));
- __global float * result = (__global float *)(res+res_idx);
- *result -= image_sum_;
+ int dst_idx = mad24(y, dst_step, mad24(x, (int)sizeof(float), dst_offset));
+ __global float * dstult = (__global float *)(dst+dst_idx);
+ *dstult -= image_sum_;
}
}
-__kernel void matchTemplate_Prepared_CCOEFF_C4 (__global const uchar * img_sums, int img_sums_step, int img_sums_offset,
- __global uchar * res, int res_step, int res_offset, int res_rows, int res_cols,
- int tpl_rows, int tpl_cols, float tpl_sum_0,float tpl_sum_1,float tpl_sum_2,float tpl_sum_3)
-{
- int gidx = get_global_id(0);
- int gidy = get_global_id(1);
-
- img_sums_step /= ELEM_SIZE;
- img_sums_offset /= ELEM_SIZE;
+#elif cn == 4
- int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float));
- float image_sum_ = 0;
+__kernel void matchTemplate_Prepared_CCOEFF(__global const uchar * src_sums, int src_sums_step, int src_sums_offset,
+ __global uchar * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols,
+ int template_rows, int template_cols, float template_sum_0, float template_sum_1, float template_sum_2, float template_sum_3)
+{
+ int x = get_global_id(0);
+ int y = get_global_id(1);
- if(gidx < res_cols && gidy < res_rows)
+ if (x < dst_cols && y < dst_rows)
{
- __global ELEM_TYPE* sum = (__global ELEM_TYPE*)(img_sums);
+ src_sums_step /= ELEM_SIZE;
+ src_sums_offset /= ELEM_SIZE;
- int c_r = SUMS_PTR(tpl_cols, tpl_rows);
- int c_o = SUMS_PTR(tpl_cols, 0);
- int o_r = SUMS_PTR(0,tpl_rows);
- int oo = SUMS_PTR(0, 0);
+ __global ELEM_TYPE* sum = (__global ELEM_TYPE*)(src_sums);
- image_sum_ += tpl_sum_0 * (float)((sum[c_r] - sum[c_o]) -(sum[o_r] - sum[oo]));
- image_sum_ += tpl_sum_1 * (float)((sum[c_r+1] - sum[c_o+1])-(sum[o_r+1] - sum[oo+1]));
- image_sum_ += tpl_sum_2 * (float)((sum[c_r+2] - sum[c_o+2])-(sum[o_r+2] - sum[oo+2]));
- image_sum_ += tpl_sum_3 * (float)((sum[c_r+3] - sum[c_o+3])-(sum[o_r+3] - sum[oo+3]));
+ int c_r = SUMS_PTR(template_cols, template_rows);
+ int c_o = SUMS_PTR(template_cols, 0);
+ int o_r = SUMS_PTR(0,template_rows);
+ int oo = SUMS_PTR(0, 0);
- __global float * result = (__global float *)(res+res_idx);
+ float image_sum_ = template_sum_0 * (float)((sum[c_r] - sum[c_o]) -(sum[o_r] - sum[oo]));
+ image_sum_ += template_sum_1 * (float)((sum[c_r+1] - sum[c_o+1])-(sum[o_r+1] - sum[oo+1]));
+ image_sum_ += template_sum_2 * (float)((sum[c_r+2] - sum[c_o+2])-(sum[o_r+2] - sum[oo+2]));
+ image_sum_ += template_sum_3 * (float)((sum[c_r+3] - sum[c_o+3])-(sum[o_r+3] - sum[oo+3]));
- *result -= image_sum_;
+ int dst_idx = mad24(y, dst_step, mad24(x, (int)sizeof(float), dst_offset));
+ __global float * dstult = (__global float *)(dst+dst_idx);
+ *dstult -= image_sum_;
}
}
-__kernel void matchTemplate_CCOEFF_NORMED_C1 (__global const uchar * img_sums, int img_sums_step, int img_sums_offset,
- __global const uchar * img_sqsums, int img_sqsums_step, int img_sqsums_offset,
- __global uchar * res, int res_step, int res_offset, int res_rows, int res_cols,
- int t_rows, int t_cols, float weight, float tpl_sum, float tpl_sqsum)
-{
- int gidx = get_global_id(0);
- int gidy = get_global_id(1);
+#else
+#error "cn should be 1, 2 or 4"
+#endif
+
+#elif defined CCOEFF_NORMED
- img_sums_offset /= ELEM_SIZE;
- img_sums_step /= ELEM_SIZE;
- img_sqsums_step /= sizeof(float);
- img_sqsums_offset /= sizeof(float);
+#if cn == 1
- int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float));
+__kernel void matchTemplate_CCOEFF_NORMED(__global const uchar * src_sums, int src_sums_step, int src_sums_offset,
+ __global const uchar * src_sqsums, int src_sqsums_step, int src_sqsums_offset,
+ __global uchar * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols,
+ int t_rows, int t_cols, float weight, float template_sum, float template_sqsum)
+{
+ int x = get_global_id(0);
+ int y = get_global_id(1);
- if(gidx < res_cols && gidy < res_rows)
+ if (x < dst_cols && y < dst_rows)
{
- __global ELEM_TYPE* sum = (__global ELEM_TYPE*)(img_sums);
- __global float * sqsum = (__global float*)(img_sqsums);
+ src_sums_offset /= ELEM_SIZE;
+ src_sums_step /= ELEM_SIZE;
+ src_sqsums_step /= sizeof(float);
+ src_sqsums_offset /= sizeof(float);
+
+ __global ELEM_TYPE* sum = (__global ELEM_TYPE*)(src_sums);
+ __global float * sqsum = (__global float*)(src_sqsums);
float image_sum_ = (float)((sum[SUMS_PTR(t_cols, t_rows)] - sum[SUMS_PTR(t_cols, 0)]) -
(sum[SUMS_PTR(0, t_rows)] - sum[SUMS_PTR(0, 0)]));
float image_sqsum_ = (float)((sqsum[SQSUMS_PTR(t_cols, t_rows)] - sqsum[SQSUMS_PTR(t_cols, 0)]) -
(sqsum[SQSUMS_PTR(0, t_rows)] - sqsum[SQSUMS_PTR(0, 0)]));
- __global float * result = (__global float *)(res+res_idx);
-
- *result = normAcc((*result) - image_sum_ * tpl_sum,
- sqrt(tpl_sqsum * (image_sqsum_ - weight * image_sum_ * image_sum_)));
+ int dst_idx = mad24(y, dst_step, mad24(x, (int)sizeof(float), dst_offset));
+ __global float * dstult = (__global float *)(dst+dst_idx);
+ *dstult = normAcc((*dstult) - image_sum_ * template_sum,
+ sqrt(template_sqsum * (image_sqsum_ - weight * image_sum_ * image_sum_)));
}
}
-__kernel void matchTemplate_CCOEFF_NORMED_C2 (__global const uchar * img_sums, int img_sums_step, int img_sums_offset,
- __global const uchar * img_sqsums, int img_sqsums_step, int img_sqsums_offset,
- __global uchar * res, int res_step, int res_offset, int res_rows, int res_cols,
- int t_rows, int t_cols, float weight, float tpl_sum_0, float tpl_sum_1, float tpl_sqsum)
-{
- int gidx = get_global_id(0);
- int gidy = get_global_id(1);
-
- img_sums_offset /= ELEM_SIZE;
- img_sums_step /= ELEM_SIZE;
- img_sqsums_step /= sizeof(float);
- img_sqsums_offset /= sizeof(float);
+#elif cn == 2
- int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float));
+__kernel void matchTemplate_CCOEFF_NORMED(__global const uchar * src_sums, int src_sums_step, int src_sums_offset,
+ __global const uchar * src_sqsums, int src_sqsums_step, int src_sqsums_offset,
+ __global uchar * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols,
+ int t_rows, int t_cols, float weight, float template_sum_0, float template_sum_1, float template_sqsum)
+{
+ int x = get_global_id(0);
+ int y = get_global_id(1);
float sum_[2];
float sqsum_[2];
- if(gidx < res_cols && gidy < res_rows)
+ if (x < dst_cols && y < dst_rows)
{
- __global ELEM_TYPE* sum = (__global ELEM_TYPE*)(img_sums);
- __global float * sqsum = (__global float*)(img_sqsums);
+ src_sums_offset /= ELEM_SIZE;
+ src_sums_step /= ELEM_SIZE;
+ src_sqsums_step /= sizeof(float);
+ src_sqsums_offset /= sizeof(float);
+
+ __global ELEM_TYPE* sum = (__global ELEM_TYPE*)(src_sums);
+ __global float * sqsum = (__global float*)(src_sqsums);
sum_[0] = (float)((sum[SUMS_PTR(t_cols, t_rows)] - sum[SUMS_PTR(t_cols, 0)])-(sum[SUMS_PTR(0, t_rows)] - sum[SUMS_PTR(0, 0)]));
sum_[1] = (float)((sum[SUMS_PTR(t_cols, t_rows)+1] - sum[SUMS_PTR(t_cols, 0)+1])-(sum[SUMS_PTR(0, t_rows)+1] - sum[SUMS_PTR(0, 0)+1]));
sqsum_[0] = (float)((sqsum[SQSUMS_PTR(t_cols, t_rows)] - sqsum[SQSUMS_PTR(t_cols, 0)])-(sqsum[SQSUMS_PTR(0, t_rows)] - sqsum[SQSUMS_PTR(0, 0)]));
sqsum_[1] = (float)((sqsum[SQSUMS_PTR(t_cols, t_rows)+1] - sqsum[SQSUMS_PTR(t_cols, 0)+1])-(sqsum[SQSUMS_PTR(0, t_rows)+1] - sqsum[SQSUMS_PTR(0, 0)+1]));
- float num = sum_[0]*tpl_sum_0 + sum_[1]*tpl_sum_1;
+ float num = sum_[0]*template_sum_0 + sum_[1]*template_sum_1;
- float denum = sqrt( tpl_sqsum * (sqsum_[0] - weight * sum_[0]* sum_[0] +
+ float denum = sqrt( template_sqsum * (sqsum_[0] - weight * sum_[0]* sum_[0] +
sqsum_[1] - weight * sum_[1]* sum_[1]));
- __global float * result = (__global float *)(res+res_idx);
- *result = normAcc((*result) - num, denum);
+ int dst_idx = mad24(y, dst_step, mad24(x, (int)sizeof(float), dst_offset));
+ __global float * dstult = (__global float *)(dst+dst_idx);
+ *dstult = normAcc((*dstult) - num, denum);
}
}
-__kernel void matchTemplate_CCOEFF_NORMED_C4 (__global const uchar * img_sums, int img_sums_step, int img_sums_offset,
- __global const uchar * img_sqsums, int img_sqsums_step, int img_sqsums_offset,
- __global uchar * res, int res_step, int res_offset, int res_rows, int res_cols,
- int t_rows, int t_cols, float weight,
- float tpl_sum_0,float tpl_sum_1,float tpl_sum_2,float tpl_sum_3,
- float tpl_sqsum)
-{
- int gidx = get_global_id(0);
- int gidy = get_global_id(1);
-
- img_sums_offset /= ELEM_SIZE;
- img_sums_step /= ELEM_SIZE;
- img_sqsums_step /= sizeof(float);
- img_sqsums_offset /= sizeof(float);
+#elif cn == 4
- int res_idx = mad24(gidy, res_step, res_offset + gidx * (int)sizeof(float));
+__kernel void matchTemplate_CCOEFF_NORMED(__global const uchar * src_sums, int src_sums_step, int src_sums_offset,
+ __global const uchar * src_sqsums, int src_sqsums_step, int src_sqsums_offset,
+ __global uchar * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols,
+ int t_rows, int t_cols, float weight,
+ float template_sum_0, float template_sum_1, float template_sum_2, float template_sum_3,
+ float template_sqsum)
+{
+ int x = get_global_id(0);
+ int y = get_global_id(1);
float sum_[4];
float sqsum_[4];
- if(gidx < res_cols && gidy < res_rows)
+ if (x < dst_cols && y < dst_rows)
{
- __global ELEM_TYPE* sum = (__global ELEM_TYPE*)(img_sums);
- __global float * sqsum = (__global float*)(img_sqsums);
+ src_sums_offset /= ELEM_SIZE;
+ src_sums_step /= ELEM_SIZE;
+ src_sqsums_step /= sizeof(float);
+ src_sqsums_offset /= sizeof(float);
+
+ __global ELEM_TYPE* sum = (__global ELEM_TYPE*)(src_sums);
+ __global float * sqsum = (__global float*)(src_sqsums);
int c_r = SUMS_PTR(t_cols, t_rows);
int c_o = SUMS_PTR(t_cols, 0);
sqsum_[2] = (float)((sqsum[c_r+2] - sqsum[c_o+2])-(sqsum[o_r+2] - sqsum[o_o+2]));
sqsum_[3] = (float)((sqsum[c_r+3] - sqsum[c_o+3])-(sqsum[o_r+3] - sqsum[o_o+3]));
- float num = sum_[0]*tpl_sum_0 + sum_[1]*tpl_sum_1 + sum_[2]*tpl_sum_2 + sum_[3]*tpl_sum_3;
+ float num = sum_[0]*template_sum_0 + sum_[1]*template_sum_1 + sum_[2]*template_sum_2 + sum_[3]*template_sum_3;
- float denum = sqrt( tpl_sqsum * (
+ float denum = sqrt( template_sqsum * (
sqsum_[0] - weight * sum_[0]* sum_[0] +
sqsum_[1] - weight * sum_[1]* sum_[1] +
sqsum_[2] - weight * sum_[2]* sum_[2] +
sqsum_[3] - weight * sum_[3]* sum_[3] ));
- __global float * result = (__global float *)(res+res_idx);
- *result = normAcc((*result) - num, denum);
+ int dst_idx = mad24(y, dst_step, mad24(x, (int)sizeof(float), dst_offset));
+ __global float * dstult = (__global float *)(dst+dst_idx);
+ *dstult = normAcc((*dstult) - num, denum);
}
}
+
+#else
+#error "cn should be 1, 2 or 4"
+#endif
+
+#endif
//M*/
#include "precomp.hpp"
+#define CV_OPENCL_RUN_ASSERT
#include "opencl_kernels.hpp"
////////////////////////////////////////////////// matchTemplate //////////////////////////////////////////////////////////
#ifdef HAVE_OPENCL
-static bool useNaive(int method, int depth, const Size & size)
+/////////////////////////////////////////////////// CCORR //////////////////////////////////////////////////////////////
+
+enum
{
-#ifdef HAVE_CLAMDFFT
- if (method == TM_SQDIFF && depth == CV_32F)
- return true;
- else if(method == TM_CCORR || (method == TM_SQDIFF && depth == CV_8U))
- return size.height < 18 && size.width < 18;
- else
+ SUM_1 = 0, SUM_2 = 1
+};
+
+static bool sumTemplate(InputArray _templ, UMat & result, int sum_type)
+{
+ CV_Assert(sum_type == SUM_1 || sum_type == SUM_2);
+
+ int type = _templ.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type);
+ int wdepth = std::max(CV_32S, depth), wtype = CV_MAKE_TYPE(wdepth, cn);
+
+ char cvt[40];
+ const char * const sumTypeToStr[] = { "SUM_1", "SUM_2" };
+ ocl::Kernel k("calcSum", ocl::imgproc::match_template_oclsrc,
+ format("-D CALC_SUM -D %s -D T=%s -D WT=%s -D convertToWT=%s -D cn=%d -D wdepth=%d",
+ sumTypeToStr[sum_type], ocl::typeToStr(type), ocl::typeToStr(wtype),
+ ocl::convertTypeStr(depth, wdepth, cn, cvt), cn, wdepth));
+ if (k.empty())
return false;
-#else
- (void)(method);
- (void)(depth);
- (void)(size);
- return true;
-#endif
-}
-/////////////////////////////////////////////////// CCORR //////////////////////////////////////////////////////////////
+ result.create(1, 1, CV_32FC1);
+ UMat templ = _templ.getUMat();
+
+ k.args(ocl::KernelArg::ReadOnly(templ), ocl::KernelArg::PtrWriteOnly(result));
+
+ return k.runTask(false);
+}
static bool matchTemplateNaive_CCORR(InputArray _image, InputArray _templ, OutputArray _result)
{
int type = _image.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type);
+ int wdepth = std::max(depth, CV_32S), wtype = CV_MAKE_TYPE(wdepth, cn);
+ char cvt[40];
ocl::Kernel k("matchTemplate_Naive_CCORR", ocl::imgproc::match_template_oclsrc,
- format("-D type=%s -D elem_type=%s -D cn=%d", ocl::typeToStr(type), ocl::typeToStr(depth), cn));
+ format("-D CCORR -D T=%s -D WT=%s -D convertToWT=%s -D cn=%d -D wdepth=%d", ocl::typeToStr(type), ocl::typeToStr(wtype),
+ ocl::convertTypeStr(depth, wdepth, cn, cvt), cn, wdepth));
if (k.empty())
return false;
UMat image = _image.getUMat(), templ = _templ.getUMat();
- _result.create(image.rows - templ.rows + 1, image.cols - templ.cols + 1, CV_32F);
+ _result.create(image.rows - templ.rows + 1, image.cols - templ.cols + 1, CV_32FC1);
UMat result = _result.getUMat();
+ k.args(ocl::KernelArg::ReadOnlyNoSize(image), ocl::KernelArg::ReadOnly(templ),
+ ocl::KernelArg::WriteOnly(result));
+
size_t globalsize[2] = { result.cols, result.rows };
- return k.args(ocl::KernelArg::ReadOnlyNoSize(image), ocl::KernelArg::ReadOnly(templ),
- ocl::KernelArg::WriteOnly(result)).run(2, globalsize, NULL, false);
+ return k.run(2, globalsize, NULL, false);
}
static bool matchTemplate_CCORR_NORMED(InputArray _image, InputArray _templ, OutputArray _result)
{
matchTemplate(_image, _templ, _result, CV_TM_CCORR);
- int type = _image.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type);
+ int type = _image.type(), cn = CV_MAT_CN(type);
ocl::Kernel k("matchTemplate_CCORR_NORMED", ocl::imgproc::match_template_oclsrc,
- format("-D type=%s -D elem_type=%s -D cn=%d", ocl::typeToStr(type),
- ocl::typeToStr(depth), cn));
+ format("-D CCORR_NORMED -D T=%s -D cn=%d", ocl::typeToStr(type), cn));
if (k.empty())
return false;
UMat image = _image.getUMat(), templ = _templ.getUMat();
- _result.create(image.rows - templ.rows + 1, image.cols - templ.cols + 1, CV_32F);
+ _result.create(image.rows - templ.rows + 1, image.cols - templ.cols + 1, CV_32FC1);
UMat result = _result.getUMat();
UMat image_sums, image_sqsums;
integral(image.reshape(1), image_sums, image_sqsums, CV_32F, CV_32F);
- UMat temp;
- multiply(templ, templ, temp, 1, CV_32F);
- Scalar s = sum(temp);
- float templ_sqsum = 0;
- for (int i = 0; i < cn; ++i)
- templ_sqsum += static_cast<float>(s[i]);
+ UMat templ_sqsum;
+ if (!sumTemplate(templ, templ_sqsum, SUM_2))
+ return false;
- size_t globalsize[2] = { result.cols, result.rows };
- return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sqsums), ocl::KernelArg::ReadWrite(result),
- templ.rows, templ.cols, templ_sqsum).run(2, globalsize, NULL, false);
-}
+ k.args(ocl::KernelArg::ReadOnlyNoSize(image_sqsums), ocl::KernelArg::ReadWrite(result),
+ templ.rows, templ.cols, ocl::KernelArg::PtrReadOnly(templ_sqsum));
-static bool matchTemplate_CCORR(InputArray _image, InputArray _templ, OutputArray _result)
-{
- if (useNaive(TM_CCORR, _image.depth(), _templ.size()) )
- return matchTemplateNaive_CCORR(_image, _templ, _result);
- else
- return false;
+ size_t globalsize[2] = { result.cols, result.rows };
+ return k.run(2, globalsize, NULL, false);
}
////////////////////////////////////// SQDIFF //////////////////////////////////////////////////////////////
static bool matchTemplateNaive_SQDIFF(InputArray _image, InputArray _templ, OutputArray _result)
{
int type = _image.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type);
+ int wdepth = std::max(depth, CV_32S), wtype = CV_MAKE_TYPE(wdepth, cn);
+ char cvt[40];
ocl::Kernel k("matchTemplate_Naive_SQDIFF", ocl::imgproc::match_template_oclsrc,
- format("-D type=%s -D elem_type=%s -D cn=%d", ocl::typeToStr(type),
- ocl::typeToStr(depth), cn));
+ format("-D SQDIFF -D T=%s -D WT=%s -D convertToWT=%s -D cn=%d -D wdepth=%d", ocl::typeToStr(type),
+ ocl::typeToStr(wtype), ocl::convertTypeStr(depth, wdepth, cn, cvt), cn, wdepth));
if (k.empty())
return false;
_result.create(image.rows - templ.rows + 1, image.cols - templ.cols + 1, CV_32F);
UMat result = _result.getUMat();
+ k.args(ocl::KernelArg::ReadOnlyNoSize(image), ocl::KernelArg::ReadOnly(templ),
+ ocl::KernelArg::WriteOnly(result));
+
size_t globalsize[2] = { result.cols, result.rows };
- return k.args(ocl::KernelArg::ReadOnlyNoSize(image), ocl::KernelArg::ReadOnly(templ),
- ocl::KernelArg::WriteOnly(result)).run(2, globalsize, NULL, false);
+ return k.run(2, globalsize, NULL, false);
}
static bool matchTemplate_SQDIFF_NORMED(InputArray _image, InputArray _templ, OutputArray _result)
{
matchTemplate(_image, _templ, _result, CV_TM_CCORR);
- int type = _image.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type);
+ int type = _image.type(), cn = CV_MAT_CN(type);
ocl::Kernel k("matchTemplate_SQDIFF_NORMED", ocl::imgproc::match_template_oclsrc,
- format("-D type=%s -D elem_type=%s -D cn=%d",
- ocl::typeToStr(type), ocl::typeToStr(depth), cn));
+ format("-D SQDIFF_NORMED -D T=%s -D cn=%d", ocl::typeToStr(type), cn));
if (k.empty())
return false;
UMat image_sums, image_sqsums;
integral(image.reshape(1), image_sums, image_sqsums, CV_32F, CV_32F);
- UMat temp;
- multiply(templ, templ, temp, 1, CV_32F);
- Scalar s = sum(temp);
- float templ_sqsum = 0;
- for (int i = 0; i < cn; ++i)
- templ_sqsum += (float)s[i];
+ UMat templ_sqsum;
+ if (!sumTemplate(_templ, templ_sqsum, SUM_2))
+ return false;
- size_t globalsize[2] = { result.cols, result.rows };
- return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sqsums), ocl::KernelArg::ReadWrite(result),
- templ.rows, templ.cols, templ_sqsum).run(2, globalsize, NULL, false);
-}
+ k.args(ocl::KernelArg::ReadOnlyNoSize(image_sqsums), ocl::KernelArg::ReadWrite(result),
+ templ.rows, templ.cols, ocl::KernelArg::PtrReadOnly(templ_sqsum));
-static bool matchTemplate_SQDIFF(InputArray _image, InputArray _templ, OutputArray _result)
-{
- if (useNaive(TM_SQDIFF, _image.depth(), _templ.size()))
- return matchTemplateNaive_SQDIFF(_image, _templ, _result);
- else
- return false;
+ size_t globalsize[2] = { result.cols, result.rows };
+ return k.run(2, globalsize, NULL, false);
}
///////////////////////////////////// CCOEFF /////////////////////////////////////////////////////////////////
UMat image_sums, temp;
integral(_image, temp);
- if(temp.depth() == CV_64F)
+ if (temp.depth() == CV_64F)
temp.convertTo(image_sums, CV_32F);
else
image_sums = temp;
int type = image_sums.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type);
- ocl::Kernel k(cv::format("matchTemplate_Prepared_CCOEFF_C%d", cn).c_str(), ocl::imgproc::match_template_oclsrc,
- format("-D type=%s -D elem_type=%s -D cn=%d", ocl::typeToStr(type), ocl::typeToStr(depth), cn));
+ ocl::Kernel k("matchTemplate_Prepared_CCOEFF", ocl::imgproc::match_template_oclsrc,
+ format("-D CCOEFF -D T=%s -D elem_type=%s -D cn=%d", ocl::typeToStr(type), ocl::typeToStr(depth), cn));
if (k.empty())
return false;
_result.create(size.height - templ.rows + 1, size.width - templ.cols + 1, CV_32F);
UMat result = _result.getUMat();
- size_t globalsize[2] = { result.cols, result.rows };
-
if (cn == 1)
{
float templ_sum = static_cast<float>(sum(_templ)[0]) / tsize.area();
- return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::ReadWrite(result),
- templ.rows, templ.cols, templ_sum).run(2, globalsize, NULL, false);
+
+ k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::ReadWrite(result),
+ templ.rows, templ.cols, templ_sum);
}
else
{
Vec4f templ_sum = Vec4f::all(0);
templ_sum = sum(templ) / tsize.area();
- if (cn == 2)
- return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::ReadWrite(result), templ.rows, templ.cols,
- templ_sum[0], templ_sum[1]).run(2, globalsize, NULL, false);
- return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::ReadWrite(result), templ.rows, templ.cols,
- templ_sum[0], templ_sum[1], templ_sum[2], templ_sum[3]).run(2, globalsize, NULL, false);
+ if (cn == 2)
+ k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::ReadWrite(result), templ.rows, templ.cols,
+ templ_sum[0], templ_sum[1]);
+ else
+ k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::ReadWrite(result), templ.rows, templ.cols,
+ templ_sum[0], templ_sum[1], templ_sum[2], templ_sum[3]);
}
+
+ size_t globalsize[2] = { result.cols, result.rows };
+ return k.run(2, globalsize, NULL, false);
}
static bool matchTemplate_CCOEFF_NORMED(InputArray _image, InputArray _templ, OutputArray _result)
int type = image_sums.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type);
- ocl::Kernel k(format("matchTemplate_CCOEFF_NORMED_C%d", cn).c_str(), ocl::imgproc::match_template_oclsrc,
- format("-D type=%s -D elem_type=%s -D cn=%d", ocl::typeToStr(type), ocl::typeToStr(depth), cn));
+ ocl::Kernel k("matchTemplate_CCOEFF_NORMED", ocl::imgproc::match_template_oclsrc,
+ format("-D CCOEFF_NORMED -D type=%s -D elem_type=%s -D cn=%d", ocl::typeToStr(type), ocl::typeToStr(depth), cn));
if (k.empty())
return false;
_result.create(size.height - templ.rows + 1, size.width - templ.cols + 1, CV_32F);
UMat result = _result.getUMat();
- size_t globalsize[2] = { result.cols, result.rows };
float scale = 1.f / tsize.area();
if (cn == 1)
return true;
}
- return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::ReadOnlyNoSize(image_sqsums),
- ocl::KernelArg::ReadWrite(result), templ.rows, templ.cols, scale, templ_sum, templ_sqsum)
- .run(2,globalsize,NULL,false);
+ k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::ReadOnlyNoSize(image_sqsums),
+ ocl::KernelArg::ReadWrite(result), templ.rows, templ.cols, scale, templ_sum, templ_sqsum);
}
else
{
}
if (cn == 2)
- return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::ReadOnlyNoSize(image_sqsums),
- ocl::KernelArg::ReadWrite(result), templ.rows, templ.cols, scale,
- templ_sum[0], templ_sum[1], templ_sqsum_sum).run(2, globalsize, NULL, false);
-
- return k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::ReadOnlyNoSize(image_sqsums),
- ocl::KernelArg::ReadWrite(result), templ.rows, templ.cols, scale,
- templ_sum[0], templ_sum[1], templ_sum[2], templ_sum[3],
- templ_sqsum_sum).run(2, globalsize, NULL, false);
+ k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::ReadOnlyNoSize(image_sqsums),
+ ocl::KernelArg::ReadWrite(result), templ.rows, templ.cols, scale,
+ templ_sum[0], templ_sum[1], templ_sqsum_sum);
+ else
+ k.args(ocl::KernelArg::ReadOnlyNoSize(image_sums), ocl::KernelArg::ReadOnlyNoSize(image_sqsums),
+ ocl::KernelArg::ReadWrite(result), templ.rows, templ.cols, scale,
+ templ_sum[0], templ_sum[1], templ_sum[2], templ_sum[3], templ_sqsum_sum);
}
+
+ size_t globalsize[2] = { result.cols, result.rows };
+ return k.run(2, globalsize, NULL, false);
}
///////////////////////////////////////////////////////////////////////////////////////////////////////////
static const Caller callers[] =
{
- matchTemplate_SQDIFF, matchTemplate_SQDIFF_NORMED, matchTemplate_CCORR,
+ matchTemplateNaive_SQDIFF, matchTemplate_SQDIFF_NORMED, matchTemplateNaive_CCORR,
matchTemplate_CCORR_NORMED, matchTemplate_CCOEFF, matchTemplate_CCOEFF_NORMED
};
const Caller caller = callers[method];
projects / platform / upstream / opencv.git / commitdiff