xref: /external/opencv3/modules/video/opencl_kernels_video.cpp

// This file is auto-generated. Do not edit!

#include "precomp.hpp"
#include "opencl_kernels_video.hpp"

namespace cv
{
namespace ocl
{
namespace video
{

const struct ProgramEntry bgfg_mog2={"bgfg_mog2",
"#if CN==1\n"
"#define T_MEAN float\n"
"#define F_ZERO (0.0f)\n"
"#define cnMode 1\n"
"#define frameToMean(a, b) (b) = *(a);\n"
"#define meanToFrame(a, b) *b = convert_uchar_sat(a);\n"
"inline float sum(float val)\n"
"{\n"
"return val;\n"
"}\n"
"#else\n"
"#define T_MEAN float4\n"
"#define F_ZERO (0.0f, 0.0f, 0.0f, 0.0f)\n"
"#define cnMode 4\n"
"#define meanToFrame(a, b)\\\n"
"b[0] = convert_uchar_sat(a.x); \\\n"
"b[1] = convert_uchar_sat(a.y); \\\n"
"b[2] = convert_uchar_sat(a.z);\n"
"#define frameToMean(a, b)\\\n"
"b.x = a[0]; \\\n"
"b.y = a[1]; \\\n"
"b.z = a[2]; \\\n"
"b.w = 0.0f;\n"
"inline float sum(const float4 val)\n"
"{\n"
"return (val.x + val.y + val.z);\n"
"}\n"
"#endif\n"
"__kernel void mog2_kernel(__global const uchar* frame, int frame_step, int frame_offset, int frame_row, int frame_col,\n"
"__global uchar* modesUsed,\n"
"__global uchar* weight,\n"
"__global uchar* mean,\n"
"__global uchar* variance,\n"
"__global uchar* fgmask, int fgmask_step, int fgmask_offset,\n"
"float alphaT, float alpha1, float prune,\n"
"float c_Tb, float c_TB, float c_Tg, float c_varMin,\n"
"float c_varMax, float c_varInit, float c_tau\n"
"#ifdef SHADOW_DETECT\n"
", uchar c_shadowVal\n"
"#endif\n"
")\n"
"{\n"
"int x = get_global_id(0);\n"
"int y = get_global_id(1);\n"
"if( x < frame_col && y < frame_row)\n"
"{\n"
"__global const uchar* _frame = (frame + mad24(y, frame_step, mad24(x, CN, frame_offset)));\n"
"T_MEAN pix;\n"
"frameToMean(_frame, pix);\n"
"uchar foreground = 255;\n"
"bool fitsPDF = false;\n"
"int pt_idx =  mad24(y, frame_col, x);\n"
"int idx_step = frame_row * frame_col;\n"
"__global uchar* _modesUsed = modesUsed + pt_idx;\n"
"uchar nmodes = _modesUsed[0];\n"
"float totalWeight = 0.0f;\n"
"__global float* _weight = (__global float*)(weight);\n"
"__global float* _variance = (__global float*)(variance);\n"
"__global T_MEAN* _mean = (__global T_MEAN*)(mean);\n"
"uchar mode = 0;\n"
"for (; mode < nmodes; ++mode)\n"
"{\n"
"int mode_idx = mad24(mode, idx_step, pt_idx);\n"
"float c_weight = mad(alpha1, _weight[mode_idx], prune);\n"
"float c_var = _variance[mode_idx];\n"
"T_MEAN c_mean = _mean[mode_idx];\n"
"T_MEAN diff = c_mean - pix;\n"
"float dist2 = dot(diff, diff);\n"
"if (totalWeight < c_TB && dist2 < c_Tb * c_var)\n"
"foreground = 0;\n"
"if (dist2 < c_Tg * c_var)\n"
"{\n"
"fitsPDF = true;\n"
"c_weight += alphaT;\n"
"float k = alphaT / c_weight;\n"
"T_MEAN mean_new = mad((T_MEAN)-k, diff, c_mean);\n"
"float variance_new  = clamp(mad(k, (dist2 - c_var), c_var), c_varMin, c_varMax);\n"
"for (int i = mode; i > 0; --i)\n"
"{\n"
"int prev_idx = mode_idx - idx_step;\n"
"if (c_weight < _weight[prev_idx])\n"
"break;\n"
"_weight[mode_idx]   = _weight[prev_idx];\n"
"_variance[mode_idx] = _variance[prev_idx];\n"
"_mean[mode_idx]     = _mean[prev_idx];\n"
"mode_idx = prev_idx;\n"
"}\n"
"_mean[mode_idx]     = mean_new;\n"
"_variance[mode_idx] = variance_new;\n"
"_weight[mode_idx]   = c_weight;\n"
"totalWeight += c_weight;\n"
"mode ++;\n"
"break;\n"
"}\n"
"if (c_weight < -prune)\n"
"c_weight = 0.0f;\n"
"_weight[mode_idx] = c_weight;\n"
"totalWeight += c_weight;\n"
"}\n"
"for (; mode < nmodes; ++mode)\n"
"{\n"
"int mode_idx = mad24(mode, idx_step, pt_idx);\n"
"float c_weight = mad(alpha1, _weight[mode_idx], prune);\n"
"if (c_weight < -prune)\n"
"{\n"
"c_weight = 0.0f;\n"
"nmodes = mode;\n"
"break;\n"
"}\n"
"_weight[mode_idx] = c_weight;\n"
"totalWeight += c_weight;\n"
"}\n"
"if (0.f < totalWeight)\n"
"{\n"
"totalWeight = 1.f / totalWeight;\n"
"for (int mode = 0; mode < nmodes; ++mode)\n"
"_weight[mad24(mode, idx_step, pt_idx)] *= totalWeight;\n"
"}\n"
"if (!fitsPDF)\n"
"{\n"
"uchar mode = nmodes == (NMIXTURES) ? (NMIXTURES) - 1 : nmodes++;\n"
"int mode_idx = mad24(mode, idx_step, pt_idx);\n"
"if (nmodes == 1)\n"
"_weight[mode_idx] = 1.f;\n"
"else\n"
"{\n"
"_weight[mode_idx] = alphaT;\n"
"for (int i = pt_idx; i < mode_idx; i += idx_step)\n"
"_weight[i] *= alpha1;\n"
"}\n"
"for (int i = nmodes - 1; i > 0; --i)\n"
"{\n"
"int prev_idx = mode_idx - idx_step;\n"
"if (alphaT < _weight[prev_idx])\n"
"break;\n"
"_weight[mode_idx]   = _weight[prev_idx];\n"
"_variance[mode_idx] = _variance[prev_idx];\n"
"_mean[mode_idx]     = _mean[prev_idx];\n"
"mode_idx = prev_idx;\n"
"}\n"
"_mean[mode_idx] = pix;\n"
"_variance[mode_idx] = c_varInit;\n"
"}\n"
"_modesUsed[0] = nmodes;\n"
"#ifdef SHADOW_DETECT\n"
"if (foreground)\n"
"{\n"
"float tWeight = 0.0f;\n"
"for (uchar mode = 0; mode < nmodes; ++mode)\n"
"{\n"
"int mode_idx = mad24(mode, idx_step, pt_idx);\n"
"T_MEAN c_mean = _mean[mode_idx];\n"
"T_MEAN pix_mean = pix * c_mean;\n"
"float numerator = sum(pix_mean);\n"
"float denominator = dot(c_mean, c_mean);\n"
"if (denominator == 0)\n"
"break;\n"
"if (numerator <= denominator && numerator >= c_tau * denominator)\n"
"{\n"
"float a = numerator / denominator;\n"
"T_MEAN dD = mad(a, c_mean, -pix);\n"
"if (dot(dD, dD) < c_Tb * _variance[mode_idx] * a * a)\n"
"{\n"
"foreground = c_shadowVal;\n"
"break;\n"
"}\n"
"}\n"
"tWeight += _weight[mode_idx];\n"
"if (tWeight > c_TB)\n"
"break;\n"
"}\n"
"}\n"
"#endif\n"
"__global uchar* _fgmask = fgmask + mad24(y, fgmask_step, x + fgmask_offset);\n"
"*_fgmask = (uchar)foreground;\n"
"}\n"
"}\n"
"__kernel void getBackgroundImage2_kernel(__global const uchar* modesUsed,\n"
"__global const uchar* weight,\n"
"__global const uchar* mean,\n"
"__global uchar* dst, int dst_step, int dst_offset, int dst_row, int dst_col,\n"
"float c_TB)\n"
"{\n"
"int x = get_global_id(0);\n"
"int y = get_global_id(1);\n"
"if(x < dst_col && y < dst_row)\n"
"{\n"
"int pt_idx =  mad24(y, dst_col, x);\n"
"__global const uchar* _modesUsed = modesUsed + pt_idx;\n"
"uchar nmodes = _modesUsed[0];\n"
"T_MEAN meanVal = (T_MEAN)F_ZERO;\n"
"float totalWeight = 0.0f;\n"
"__global const float* _weight = (__global const float*)weight;\n"
"__global const T_MEAN* _mean = (__global const T_MEAN*)(mean);\n"
"int idx_step = dst_row * dst_col;\n"
"for (uchar mode = 0; mode < nmodes; ++mode)\n"
"{\n"
"int mode_idx = mad24(mode, idx_step, pt_idx);\n"
"float c_weight = _weight[mode_idx];\n"
"T_MEAN c_mean = _mean[mode_idx];\n"
"meanVal = mad(c_weight, c_mean, meanVal);\n"
"totalWeight += c_weight;\n"
"if (totalWeight > c_TB)\n"
"break;\n"
"}\n"
"if (0.f < totalWeight)\n"
"meanVal = meanVal / totalWeight;\n"
"else\n"
"meanVal = (T_MEAN)(0.f);\n"
"__global uchar* _dst = dst + mad24(y, dst_step, mad24(x, CN, dst_offset));\n"
"meanToFrame(meanVal, _dst);\n"
"}\n"
"}\n"
, "b6e3850899862b7f0ab67cb32f1d52e9"};
ProgramSource bgfg_mog2_oclsrc(bgfg_mog2.programStr);
const struct ProgramEntry optical_flow_farneback={"optical_flow_farneback",
"#define tx  (int)get_local_id(0)\n"
"#define ty  get_local_id(1)\n"
"#define bx  get_group_id(0)\n"
"#define bdx (int)get_local_size(0)\n"
"#define BORDER_SIZE 5\n"
"#define MAX_KSIZE_HALF 100\n"
"#ifndef polyN\n"
"#define polyN 5\n"
"#endif\n"
"#if USE_DOUBLE\n"
"#ifdef cl_amd_fp64\n"
"#pragma OPENCL EXTENSION cl_amd_fp64:enable\n"
"#elif defined (cl_khr_fp64)\n"
"#pragma OPENCL EXTENSION cl_khr_fp64:enable\n"
"#endif\n"
"#define TYPE double\n"
"#define VECTYPE double4\n"
"#else\n"
"#define TYPE float\n"
"#define VECTYPE float4\n"
"#endif\n"
"__kernel void polynomialExpansion(__global __const float * src, int srcStep,\n"
"__global float * dst, int dstStep,\n"
"const int rows, const  int cols,\n"
"__global __const float * c_g,\n"
"__global __const float * c_xg,\n"
"__global __const float * c_xxg,\n"
"__local float * smem,\n"
"const VECTYPE ig)\n"
"{\n"
"const int y = get_global_id(1);\n"
"const int x = bx * (bdx - 2*polyN) + tx - polyN;\n"
"int xWarped;\n"
"__local float *row = smem + tx;\n"
"if (y < rows && y >= 0)\n"
"{\n"
"xWarped = min(max(x, 0), cols - 1);\n"
"row[0] = src[mad24(y, srcStep, xWarped)] * c_g[0];\n"
"row[bdx] = 0.f;\n"
"row[2*bdx] = 0.f;\n"
"#pragma unroll\n"
"for (int k = 1; k <= polyN; ++k)\n"
"{\n"
"float t0 = src[mad24(max(y - k, 0), srcStep, xWarped)];\n"
"float t1 = src[mad24(min(y + k, rows - 1), srcStep, xWarped)];\n"
"row[0] += c_g[k] * (t0 + t1);\n"
"row[bdx] += c_xg[k] * (t1 - t0);\n"
"row[2*bdx] += c_xxg[k] * (t0 + t1);\n"
"}\n"
"}\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"if (y < rows && y >= 0 && tx >= polyN && tx + polyN < bdx && x < cols)\n"
"{\n"
"TYPE b1 = c_g[0] * row[0];\n"
"TYPE b3 = c_g[0] * row[bdx];\n"
"TYPE b5 = c_g[0] * row[2*bdx];\n"
"TYPE b2 = 0, b4 = 0, b6 = 0;\n"
"#pragma unroll\n"
"for (int k = 1; k <= polyN; ++k)\n"
"{\n"
"b1 += (row[k] + row[-k]) * c_g[k];\n"
"b4 += (row[k] + row[-k]) * c_xxg[k];\n"
"b2 += (row[k] - row[-k]) * c_xg[k];\n"
"b3 += (row[k + bdx] + row[-k + bdx]) * c_g[k];\n"
"b6 += (row[k + bdx] - row[-k + bdx]) * c_xg[k];\n"
"b5 += (row[k + 2*bdx] + row[-k + 2*bdx]) * c_g[k];\n"
"}\n"
"dst[mad24(y, dstStep, xWarped)] = (float)(b3*ig.s0);\n"
"dst[mad24(rows + y, dstStep, xWarped)] = (float)(b2*ig.s0);\n"
"dst[mad24(2*rows + y, dstStep, xWarped)] = (float)(b1*ig.s1 + b5*ig.s2);\n"
"dst[mad24(3*rows + y, dstStep, xWarped)] = (float)(b1*ig.s1 + b4*ig.s2);\n"
"dst[mad24(4*rows + y, dstStep, xWarped)] = (float)(b6*ig.s3);\n"
"}\n"
"}\n"
"inline int idx_row_low(const int y, const int last_row)\n"
"{\n"
"return abs(y) % (last_row + 1);\n"
"}\n"
"inline int idx_row_high(const int y, const int last_row)\n"
"{\n"
"return abs(last_row - abs(last_row - y)) % (last_row + 1);\n"
"}\n"
"inline int idx_col_low(const int x, const int last_col)\n"
"{\n"
"return abs(x) % (last_col + 1);\n"
"}\n"
"inline int idx_col_high(const int x, const int last_col)\n"
"{\n"
"return abs(last_col - abs(last_col - x)) % (last_col + 1);\n"
"}\n"
"inline int idx_col(const int x, const int last_col)\n"
"{\n"
"return idx_col_low(idx_col_high(x, last_col), last_col);\n"
"}\n"
"__kernel void gaussianBlur(__global const float * src, int srcStep,\n"
"__global float * dst, int dstStep, const int rows, const  int cols,\n"
"__global const float * c_gKer, const int ksizeHalf,\n"
"__local float * smem)\n"
"{\n"
"const int y = get_global_id(1);\n"
"const int x = get_global_id(0);\n"
"__local float *row = smem + ty * (bdx + 2*ksizeHalf);\n"
"if (y < rows)\n"
"{\n"
"for (int i = tx; i < bdx + 2*ksizeHalf; i += bdx)\n"
"{\n"
"int xExt = (int)(bx * bdx) + i - ksizeHalf;\n"
"xExt = idx_col(xExt, cols - 1);\n"
"row[i] = src[mad24(y, srcStep, xExt)] * c_gKer[0];\n"
"for (int j = 1; j <= ksizeHalf; ++j)\n"
"row[i] += (src[mad24(idx_row_low(y - j, rows - 1), srcStep, xExt)]\n"
"+ src[mad24(idx_row_high(y + j, rows - 1), srcStep, xExt)]) * c_gKer[j];\n"
"}\n"
"}\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"if (y < rows && y >= 0 && x < cols && x >= 0)\n"
"{\n"
"row += tx + ksizeHalf;\n"
"float res = row[0] * c_gKer[0];\n"
"for (int i = 1; i <= ksizeHalf; ++i)\n"
"res += (row[-i] + row[i]) * c_gKer[i];\n"
"dst[mad24(y, dstStep, x)] = res;\n"
"}\n"
"}\n"
"__kernel void gaussianBlur5(__global const float * src, int srcStep,\n"
"__global float * dst, int dstStep,\n"
"const int rows, const  int cols,\n"
"__global const float * c_gKer, const int ksizeHalf,\n"
"__local float * smem)\n"
"{\n"
"const int y = get_global_id(1);\n"
"const int x = get_global_id(0);\n"
"const int smw = bdx + 2*ksizeHalf;\n"
"__local volatile float *row = smem + 5 * ty * smw;\n"
"if (y < rows)\n"
"{\n"
"for (int i = tx; i < bdx + 2*ksizeHalf; i += bdx)\n"
"{\n"
"int xExt = (int)(bx * bdx) + i - ksizeHalf;\n"
"xExt = idx_col(xExt, cols - 1);\n"
"#pragma unroll\n"
"for (int k = 0; k < 5; ++k)\n"
"row[k*smw + i] = src[mad24(k*rows + y, srcStep, xExt)] * c_gKer[0];\n"
"for (int j = 1; j <= ksizeHalf; ++j)\n"
"#pragma unroll\n"
"for (int k = 0; k < 5; ++k)\n"
"row[k*smw + i] +=\n"
"(src[mad24(k*rows + idx_row_low(y - j, rows - 1), srcStep, xExt)] +\n"
"src[mad24(k*rows + idx_row_high(y + j, rows - 1), srcStep, xExt)]) * c_gKer[j];\n"
"}\n"
"}\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"if (y < rows && y >= 0 && x < cols && x >= 0)\n"
"{\n"
"row += tx + ksizeHalf;\n"
"float res[5];\n"
"#pragma unroll\n"
"for (int k = 0; k < 5; ++k)\n"
"res[k] = row[k*smw] * c_gKer[0];\n"
"for (int i = 1; i <= ksizeHalf; ++i)\n"
"#pragma unroll\n"
"for (int k = 0; k < 5; ++k)\n"
"res[k] += (row[k*smw - i] + row[k*smw + i]) * c_gKer[i];\n"
"#pragma unroll\n"
"for (int k = 0; k < 5; ++k)\n"
"dst[mad24(k*rows + y, dstStep, x)] = res[k];\n"
"}\n"
"}\n"
"__constant float c_border[BORDER_SIZE + 1] = { 0.14f, 0.14f, 0.4472f, 0.4472f, 0.4472f, 1.f };\n"
"__kernel void updateMatrices(__global const float * flowx, int xStep,\n"
"__global const float * flowy, int yStep,\n"
"const int rows, const int cols,\n"
"__global const float * R0, int R0Step,\n"
"__global const float * R1, int R1Step,\n"
"__global float * M, int mStep)\n"
"{\n"
"const int y = get_global_id(1);\n"
"const int x = get_global_id(0);\n"
"if (y < rows && y >= 0 && x < cols && x >= 0)\n"
"{\n"
"float dx = flowx[mad24(y, xStep, x)];\n"
"float dy = flowy[mad24(y, yStep, x)];\n"
"float fx = x + dx;\n"
"float fy = y + dy;\n"
"int x1 = convert_int(floor(fx));\n"
"int y1 = convert_int(floor(fy));\n"
"fx -= x1;\n"
"fy -= y1;\n"
"float r2, r3, r4, r5, r6;\n"
"if (x1 >= 0 && y1 >= 0 && x1 < cols - 1 && y1 < rows - 1)\n"
"{\n"
"float a00 = (1.f - fx) * (1.f - fy);\n"
"float a01 = fx * (1.f - fy);\n"
"float a10 = (1.f - fx) * fy;\n"
"float a11 = fx * fy;\n"
"r2 = a00 * R1[mad24(y1, R1Step, x1)] +\n"
"a01 * R1[mad24(y1, R1Step, x1 + 1)] +\n"
"a10 * R1[mad24(y1 + 1, R1Step, x1)] +\n"
"a11 * R1[mad24(y1 + 1, R1Step, x1 + 1)];\n"
"r3 = a00 * R1[mad24(rows + y1, R1Step, x1)] +\n"
"a01 * R1[mad24(rows + y1, R1Step, x1 + 1)] +\n"
"a10 * R1[mad24(rows + y1 + 1, R1Step, x1)] +\n"
"a11 * R1[mad24(rows + y1 + 1, R1Step, x1 + 1)];\n"
"r4 = a00 * R1[mad24(2*rows + y1, R1Step, x1)] +\n"
"a01 * R1[mad24(2*rows + y1, R1Step, x1 + 1)] +\n"
"a10 * R1[mad24(2*rows + y1 + 1, R1Step, x1)] +\n"
"a11 * R1[mad24(2*rows + y1 + 1, R1Step, x1 + 1)];\n"
"r5 = a00 * R1[mad24(3*rows + y1, R1Step, x1)] +\n"
"a01 * R1[mad24(3*rows + y1, R1Step, x1 + 1)] +\n"
"a10 * R1[mad24(3*rows + y1 + 1, R1Step, x1)] +\n"
"a11 * R1[mad24(3*rows + y1 + 1, R1Step, x1 + 1)];\n"
"r6 = a00 * R1[mad24(4*rows + y1, R1Step, x1)] +\n"
"a01 * R1[mad24(4*rows + y1, R1Step, x1 + 1)] +\n"
"a10 * R1[mad24(4*rows + y1 + 1, R1Step, x1)] +\n"
"a11 * R1[mad24(4*rows + y1 + 1, R1Step, x1 + 1)];\n"
"r4 = (R0[mad24(2*rows + y, R0Step, x)] + r4) * 0.5f;\n"
"r5 = (R0[mad24(3*rows + y, R0Step, x)] + r5) * 0.5f;\n"
"r6 = (R0[mad24(4*rows + y, R0Step, x)] + r6) * 0.25f;\n"
"}\n"
"else\n"
"{\n"
"r2 = r3 = 0.f;\n"
"r4 = R0[mad24(2*rows + y, R0Step, x)];\n"
"r5 = R0[mad24(3*rows + y, R0Step, x)];\n"
"r6 = R0[mad24(4*rows + y, R0Step, x)] * 0.5f;\n"
"}\n"
"r2 = (R0[mad24(y, R0Step, x)] - r2) * 0.5f;\n"
"r3 = (R0[mad24(rows + y, R0Step, x)] - r3) * 0.5f;\n"
"r2 += r4*dy + r6*dx;\n"
"r3 += r6*dy + r5*dx;\n"
"float scale =\n"
"c_border[min(x, BORDER_SIZE)] *\n"
"c_border[min(y, BORDER_SIZE)] *\n"
"c_border[min(cols - x - 1, BORDER_SIZE)] *\n"
"c_border[min(rows - y - 1, BORDER_SIZE)];\n"
"r2 *= scale;\n"
"r3 *= scale;\n"
"r4 *= scale;\n"
"r5 *= scale;\n"
"r6 *= scale;\n"
"M[mad24(y, mStep, x)] = r4*r4 + r6*r6;\n"
"M[mad24(rows + y, mStep, x)] = (r4 + r5)*r6;\n"
"M[mad24(2*rows + y, mStep, x)] = r5*r5 + r6*r6;\n"
"M[mad24(3*rows + y, mStep, x)] = r4*r2 + r6*r3;\n"
"M[mad24(4*rows + y, mStep, x)] = r6*r2 + r5*r3;\n"
"}\n"
"}\n"
"__kernel void boxFilter5(__global const float * src, int srcStep,\n"
"__global float * dst, int dstStep,\n"
"const int rows, const  int cols,\n"
"const int ksizeHalf,\n"
"__local float * smem)\n"
"{\n"
"const int y = get_global_id(1);\n"
"const int x = get_global_id(0);\n"
"const float boxAreaInv = 1.f / ((1 + 2*ksizeHalf) * (1 + 2*ksizeHalf));\n"
"const int smw = bdx + 2*ksizeHalf;\n"
"__local float *row = smem + 5 * ty * smw;\n"
"if (y < rows)\n"
"{\n"
"for (int i = tx; i < bdx + 2*ksizeHalf; i += bdx)\n"
"{\n"
"int xExt = (int)(bx * bdx) + i - ksizeHalf;\n"
"xExt = min(max(xExt, 0), cols - 1);\n"
"#pragma unroll\n"
"for (int k = 0; k < 5; ++k)\n"
"row[k*smw + i] = src[mad24(k*rows + y, srcStep, xExt)];\n"
"for (int j = 1; j <= ksizeHalf; ++j)\n"
"#pragma unroll\n"
"for (int k = 0; k < 5; ++k)\n"
"row[k*smw + i] +=\n"
"src[mad24(k*rows + max(y - j, 0), srcStep, xExt)] +\n"
"src[mad24(k*rows + min(y + j, rows - 1), srcStep, xExt)];\n"
"}\n"
"}\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"if (y < rows && y >= 0 && x < cols && x >= 0)\n"
"{\n"
"row += tx + ksizeHalf;\n"
"float res[5];\n"
"#pragma unroll\n"
"for (int k = 0; k < 5; ++k)\n"
"res[k] = row[k*smw];\n"
"for (int i = 1; i <= ksizeHalf; ++i)\n"
"#pragma unroll\n"
"for (int k = 0; k < 5; ++k)\n"
"res[k] += row[k*smw - i] + row[k*smw + i];\n"
"#pragma unroll\n"
"for (int k = 0; k < 5; ++k)\n"
"dst[mad24(k*rows + y, dstStep, x)] = res[k] * boxAreaInv;\n"
"}\n"
"}\n"
"__kernel void updateFlow(__global const float * M, int mStep,\n"
"__global float * flowx, int xStep,\n"
"__global float * flowy, int yStep,\n"
"const int rows, const int cols)\n"
"{\n"
"const int y = get_global_id(1);\n"
"const int x = get_global_id(0);\n"
"if (y < rows && y >= 0 && x < cols && x >= 0)\n"
"{\n"
"float g11 = M[mad24(y, mStep, x)];\n"
"float g12 = M[mad24(rows + y, mStep, x)];\n"
"float g22 = M[mad24(2*rows + y, mStep, x)];\n"
"float h1 =  M[mad24(3*rows + y, mStep, x)];\n"
"float h2 =  M[mad24(4*rows + y, mStep, x)];\n"
"float detInv = 1.f / (g11*g22 - g12*g12 + 1e-3f);\n"
"flowx[mad24(y, xStep, x)] = (g11*h2 - g12*h1) * detInv;\n"
"flowy[mad24(y, yStep, x)] = (g22*h1 - g12*h2) * detInv;\n"
"}\n"
"}\n"
, "529300e6242f574f83d11a089cc120c0"};
ProgramSource optical_flow_farneback_oclsrc(optical_flow_farneback.programStr);
const struct ProgramEntry optical_flow_tvl1={"optical_flow_tvl1",
"__kernel void centeredGradientKernel(__global const float* src_ptr, int src_col, int src_row, int src_step,\n"
"__global float* dx, __global float* dy, int d_step)\n"
"{\n"
"int x = get_global_id(0);\n"
"int y = get_global_id(1);\n"
"if((x < src_col)&&(y < src_row))\n"
"{\n"
"int src_x1 = (x + 1) < (src_col -1)? (x + 1) : (src_col - 1);\n"
"int src_x2 = (x - 1) > 0 ? (x -1) : 0;\n"
"dx[y * d_step+ x] = 0.5f * (src_ptr[y * src_step + src_x1] - src_ptr[y * src_step+ src_x2]);\n"
"int src_y1 = (y+1) < (src_row - 1) ? (y + 1) : (src_row - 1);\n"
"int src_y2 = (y - 1) > 0 ? (y - 1) : 0;\n"
"dy[y * d_step+ x] = 0.5f * (src_ptr[src_y1 * src_step + x] - src_ptr[src_y2 * src_step+ x]);\n"
"}\n"
"}\n"
"inline float bicubicCoeff(float x_)\n"
"{\n"
"float x = fabs(x_);\n"
"if (x <= 1.0f)\n"
"return x * x * (1.5f * x - 2.5f) + 1.0f;\n"
"else if (x < 2.0f)\n"
"return x * (x * (-0.5f * x + 2.5f) - 4.0f) + 2.0f;\n"
"else\n"
"return 0.0f;\n"
"}\n"
"__kernel void warpBackwardKernel(__global const float* I0, int I0_step, int I0_col, int I0_row,\n"
"image2d_t tex_I1, image2d_t tex_I1x, image2d_t tex_I1y,\n"
"__global const float* u1, int u1_step,\n"
"__global const float* u2,\n"
"__global float* I1w,\n"
"__global float* I1wx, \n"
"__global float* I1wy, \n"
"__global float* grad, \n"
"__global float* rho,\n"
"int I1w_step,\n"
"int u2_step,\n"
"int u1_offset_x,\n"
"int u1_offset_y,\n"
"int u2_offset_x,\n"
"int u2_offset_y)\n"
"{\n"
"int x = get_global_id(0);\n"
"int y = get_global_id(1);\n"
"if(x < I0_col&&y < I0_row)\n"
"{\n"
"float u1Val = u1[(y + u1_offset_y) * u1_step + x + u1_offset_x];\n"
"float u2Val = u2[(y + u2_offset_y) * u2_step + x + u2_offset_x];\n"
"float wx = x + u1Val;\n"
"float wy = y + u2Val;\n"
"int xmin = ceil(wx - 2.0f);\n"
"int xmax = floor(wx + 2.0f);\n"
"int ymin = ceil(wy - 2.0f);\n"
"int ymax = floor(wy + 2.0f);\n"
"float sum  = 0.0f;\n"
"float sumx = 0.0f;\n"
"float sumy = 0.0f;\n"
"float wsum = 0.0f;\n"
"sampler_t sampleri = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;\n"
"for (int cy = ymin; cy <= ymax; ++cy)\n"
"{\n"
"for (int cx = xmin; cx <= xmax; ++cx)\n"
"{\n"
"float w = bicubicCoeff(wx - cx) * bicubicCoeff(wy - cy);\n"
"int2 cood = (int2)(cx, cy);\n"
"sum += w * read_imagef(tex_I1, sampleri, cood).x;\n"
"sumx += w * read_imagef(tex_I1x, sampleri, cood).x;\n"
"sumy += w * read_imagef(tex_I1y, sampleri, cood).x;\n"
"wsum += w;\n"
"}\n"
"}\n"
"float coeff = 1.0f / wsum;\n"
"float I1wVal  = sum  * coeff;\n"
"float I1wxVal = sumx * coeff;\n"
"float I1wyVal = sumy * coeff;\n"
"I1w[y * I1w_step + x]  = I1wVal;\n"
"I1wx[y * I1w_step + x] = I1wxVal;\n"
"I1wy[y * I1w_step + x] = I1wyVal;\n"
"float Ix2 = I1wxVal * I1wxVal;\n"
"float Iy2 = I1wyVal * I1wyVal;\n"
"grad[y * I1w_step + x] = Ix2 + Iy2;\n"
"float I0Val = I0[y * I0_step + x];\n"
"rho[y * I1w_step + x] = I1wVal - I1wxVal * u1Val - I1wyVal * u2Val - I0Val;\n"
"}\n"
"}\n"
"inline float readImage(__global float *image,  int x,  int y,  int rows,  int cols, int elemCntPerRow)\n"
"{\n"
"int i0 = clamp(x, 0, cols - 1);\n"
"int j0 = clamp(y, 0, rows - 1);\n"
"return image[j0 * elemCntPerRow + i0];\n"
"}\n"
"__kernel void warpBackwardKernelNoImage2d(__global const float* I0, int I0_step, int I0_col, int I0_row,\n"
"__global const float* tex_I1, __global const float* tex_I1x, __global const float* tex_I1y,\n"
"__global const float* u1, int u1_step,\n"
"__global const float* u2,\n"
"__global float* I1w,\n"
"__global float* I1wx, \n"
"__global float* I1wy, \n"
"__global float* grad, \n"
"__global float* rho,\n"
"int I1w_step,\n"
"int u2_step,\n"
"int I1_step,\n"
"int I1x_step)\n"
"{\n"
"int x = get_global_id(0);\n"
"int y = get_global_id(1);\n"
"if(x < I0_col&&y < I0_row)\n"
"{\n"
"float u1Val = u1[y * u1_step + x];\n"
"float u2Val = u2[y * u2_step + x];\n"
"float wx = x + u1Val;\n"
"float wy = y + u2Val;\n"
"int xmin = ceil(wx - 2.0f);\n"
"int xmax = floor(wx + 2.0f);\n"
"int ymin = ceil(wy - 2.0f);\n"
"int ymax = floor(wy + 2.0f);\n"
"float sum  = 0.0f;\n"
"float sumx = 0.0f;\n"
"float sumy = 0.0f;\n"
"float wsum = 0.0f;\n"
"for (int cy = ymin; cy <= ymax; ++cy)\n"
"{\n"
"for (int cx = xmin; cx <= xmax; ++cx)\n"
"{\n"
"float w = bicubicCoeff(wx - cx) * bicubicCoeff(wy - cy);\n"
"int2 cood = (int2)(cx, cy);\n"
"sum += w * readImage(tex_I1, cood.x, cood.y, I0_col, I0_row, I1_step);\n"
"sumx += w * readImage(tex_I1x, cood.x, cood.y, I0_col, I0_row, I1x_step);\n"
"sumy += w * readImage(tex_I1y, cood.x, cood.y, I0_col, I0_row, I1x_step);\n"
"wsum += w;\n"
"}\n"
"}\n"
"float coeff = 1.0f / wsum;\n"
"float I1wVal  = sum  * coeff;\n"
"float I1wxVal = sumx * coeff;\n"
"float I1wyVal = sumy * coeff;\n"
"I1w[y * I1w_step + x]  = I1wVal;\n"
"I1wx[y * I1w_step + x] = I1wxVal;\n"
"I1wy[y * I1w_step + x] = I1wyVal;\n"
"float Ix2 = I1wxVal * I1wxVal;\n"
"float Iy2 = I1wyVal * I1wyVal;\n"
"grad[y * I1w_step + x] = Ix2 + Iy2;\n"
"float I0Val = I0[y * I0_step + x];\n"
"rho[y * I1w_step + x] = I1wVal - I1wxVal * u1Val - I1wyVal * u2Val - I0Val;\n"
"}\n"
"}\n"
"__kernel void estimateDualVariablesKernel(__global const float* u1, int u1_col, int u1_row, int u1_step,\n"
"__global const float* u2,\n"
"__global float* p11, int p11_step,\n"
"__global float* p12,\n"
"__global float* p21,\n"
"__global float* p22,\n"
"float taut,\n"
"int u2_step,\n"
"int u1_offset_x,\n"
"int u1_offset_y,\n"
"int u2_offset_x,\n"
"int u2_offset_y)\n"
"{\n"
"int x = get_global_id(0);\n"
"int y = get_global_id(1);\n"
"if(x < u1_col && y < u1_row)\n"
"{\n"
"int src_x1 = (x + 1) < (u1_col - 1) ? (x + 1) : (u1_col - 1);\n"
"float u1x = u1[(y + u1_offset_y) * u1_step + src_x1 + u1_offset_x] - u1[(y + u1_offset_y) * u1_step + x + u1_offset_x];\n"
"int src_y1 = (y + 1) < (u1_row - 1) ? (y + 1) : (u1_row - 1);\n"
"float u1y = u1[(src_y1 + u1_offset_y) * u1_step + x + u1_offset_x] - u1[(y + u1_offset_y) * u1_step + x + u1_offset_x];\n"
"int src_x2 = (x + 1) < (u1_col - 1) ? (x + 1) : (u1_col - 1);\n"
"float u2x = u2[(y + u2_offset_y) * u2_step + src_x2 + u2_offset_x] - u2[(y + u2_offset_y) * u2_step + x + u2_offset_x];\n"
"int src_y2 = (y + 1) <  (u1_row - 1) ? (y + 1) : (u1_row - 1);\n"
"float u2y = u2[(src_y2 + u2_offset_y) * u2_step + x + u2_offset_x] - u2[(y + u2_offset_y) * u2_step + x + u2_offset_x];\n"
"float g1 = hypot(u1x, u1y);\n"
"float g2 = hypot(u2x, u2y);\n"
"float ng1 = 1.0f + taut * g1;\n"
"float ng2 = 1.0f + taut * g2;\n"
"p11[y * p11_step + x] = (p11[y * p11_step + x] + taut * u1x) / ng1;\n"
"p12[y * p11_step + x] = (p12[y * p11_step + x] + taut * u1y) / ng1;\n"
"p21[y * p11_step + x] = (p21[y * p11_step + x] + taut * u2x) / ng2;\n"
"p22[y * p11_step + x] = (p22[y * p11_step + x] + taut * u2y) / ng2;\n"
"}\n"
"}\n"
"inline float divergence(__global const float* v1, __global const float* v2, int y, int x, int v1_step, int v2_step)\n"
"{\n"
"if (x > 0 && y > 0)\n"
"{\n"
"float v1x = v1[y * v1_step + x] - v1[y * v1_step + x - 1];\n"
"float v2y = v2[y * v2_step + x] - v2[(y - 1) * v2_step + x];\n"
"return v1x + v2y;\n"
"}\n"
"else\n"
"{\n"
"if (y > 0)\n"
"return v1[y * v1_step + 0] + v2[y * v2_step + 0] - v2[(y - 1) * v2_step + 0];\n"
"else\n"
"{\n"
"if (x > 0)\n"
"return v1[0 * v1_step + x] - v1[0 * v1_step + x - 1] + v2[0 * v2_step + x];\n"
"else\n"
"return v1[0 * v1_step + 0] + v2[0 * v2_step + 0];\n"
"}\n"
"}\n"
"}\n"
"__kernel void estimateUKernel(__global const float* I1wx, int I1wx_col, int I1wx_row, int I1wx_step,\n"
"__global const float* I1wy, \n"
"__global const float* grad, \n"
"__global const float* rho_c, \n"
"__global const float* p11, \n"
"__global const float* p12, \n"
"__global const float* p21, \n"
"__global const float* p22, \n"
"__global float* u1, int u1_step,\n"
"__global float* u2,\n"
"__global float* error, float l_t, float theta, int u2_step,\n"
"int u1_offset_x,\n"
"int u1_offset_y,\n"
"int u2_offset_x,\n"
"int u2_offset_y,\n"
"char calc_error)\n"
"{\n"
"int x = get_global_id(0);\n"
"int y = get_global_id(1);\n"
"if(x < I1wx_col && y < I1wx_row)\n"
"{\n"
"float I1wxVal = I1wx[y * I1wx_step + x];\n"
"float I1wyVal = I1wy[y * I1wx_step + x];\n"
"float gradVal = grad[y * I1wx_step + x];\n"
"float u1OldVal = u1[(y + u1_offset_y) * u1_step + x + u1_offset_x];\n"
"float u2OldVal = u2[(y + u2_offset_y) * u2_step + x + u2_offset_x];\n"
"float rho = rho_c[y * I1wx_step + x] + (I1wxVal * u1OldVal + I1wyVal * u2OldVal);\n"
"float d1 = 0.0f;\n"
"float d2 = 0.0f;\n"
"if (rho < -l_t * gradVal)\n"
"{\n"
"d1 = l_t * I1wxVal;\n"
"d2 = l_t * I1wyVal;\n"
"}\n"
"else if (rho > l_t * gradVal)\n"
"{\n"
"d1 = -l_t * I1wxVal;\n"
"d2 = -l_t * I1wyVal;\n"
"}\n"
"else if (gradVal > 1.192092896e-07f)\n"
"{\n"
"float fi = -rho / gradVal;\n"
"d1 = fi * I1wxVal;\n"
"d2 = fi * I1wyVal;\n"
"}\n"
"float v1 = u1OldVal + d1;\n"
"float v2 = u2OldVal + d2;\n"
"float div_p1 = divergence(p11, p12, y, x, I1wx_step, I1wx_step);\n"
"float div_p2 = divergence(p21, p22, y, x, I1wx_step, I1wx_step);\n"
"float u1NewVal = v1 + theta * div_p1;\n"
"float u2NewVal = v2 + theta * div_p2;\n"
"u1[(y + u1_offset_y) * u1_step + x + u1_offset_x] = u1NewVal;\n"
"u2[(y + u2_offset_y) * u2_step + x + u2_offset_x] = u2NewVal;\n"
"if(calc_error)\n"
"{\n"
"float n1 = (u1OldVal - u1NewVal) * (u1OldVal - u1NewVal);\n"
"float n2 = (u2OldVal - u2NewVal) * (u2OldVal - u2NewVal);\n"
"error[y * I1wx_step + x] = n1 + n2;\n"
"}\n"
"}\n"
"}\n"
, "a9d306a49b405703820fae23312ebd28"};
ProgramSource optical_flow_tvl1_oclsrc(optical_flow_tvl1.programStr);
const struct ProgramEntry pyrlk={"pyrlk",
"#define GRIDSIZE    3\n"
"#define LSx 8\n"
"#define LSy 8\n"
"#define LM_W (LSx*GRIDSIZE+2)\n"
"#define LM_H (LSy*GRIDSIZE+2)\n"
"#define BUFFER  (LSx*LSy)\n"
"#define BUFFER2 BUFFER>>1\n"
"#ifndef WAVE_SIZE\n"
"#define WAVE_SIZE 1\n"
"#endif\n"
"#ifdef CPU\n"
"inline void reduce3(float val1, float val2, float val3,  __local float* smem1,  __local float* smem2,  __local float* smem3, int tid)\n"
"{\n"
"smem1[tid] = val1;\n"
"smem2[tid] = val2;\n"
"smem3[tid] = val3;\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"for(int i = BUFFER2; i > 0; i >>= 1)\n"
"{\n"
"if(tid < i)\n"
"{\n"
"smem1[tid] += smem1[tid + i];\n"
"smem2[tid] += smem2[tid + i];\n"
"smem3[tid] += smem3[tid + i];\n"
"}\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"}\n"
"}\n"
"inline void reduce2(float val1, float val2, volatile __local float* smem1, volatile __local float* smem2, int tid)\n"
"{\n"
"smem1[tid] = val1;\n"
"smem2[tid] = val2;\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"for(int i = BUFFER2; i > 0; i >>= 1)\n"
"{\n"
"if(tid < i)\n"
"{\n"
"smem1[tid] += smem1[tid + i];\n"
"smem2[tid] += smem2[tid + i];\n"
"}\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"}\n"
"}\n"
"inline void reduce1(float val1, volatile __local float* smem1, int tid)\n"
"{\n"
"smem1[tid] = val1;\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"for(int i = BUFFER2; i > 0; i >>= 1)\n"
"{\n"
"if(tid < i)\n"
"{\n"
"smem1[tid] += smem1[tid + i];\n"
"}\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"}\n"
"}\n"
"#else\n"
"inline void reduce3(float val1, float val2, float val3,\n"
"__local volatile float* smem1, __local volatile float* smem2, __local volatile float* smem3, int tid)\n"
"{\n"
"smem1[tid] = val1;\n"
"smem2[tid] = val2;\n"
"smem3[tid] = val3;\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"if (tid < 32)\n"
"{\n"
"smem1[tid] += smem1[tid + 32];\n"
"smem2[tid] += smem2[tid + 32];\n"
"smem3[tid] += smem3[tid + 32];\n"
"#if WAVE_SIZE < 32\n"
"}\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"if (tid < 16)\n"
"{\n"
"#endif\n"
"smem1[tid] += smem1[tid + 16];\n"
"smem2[tid] += smem2[tid + 16];\n"
"smem3[tid] += smem3[tid + 16];\n"
"#if WAVE_SIZE <16\n"
"}\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"if (tid<1)\n"
"{\n"
"#endif\n"
"local float8* m1 = (local float8*)smem1;\n"
"local float8* m2 = (local float8*)smem2;\n"
"local float8* m3 = (local float8*)smem3;\n"
"float8 t1 = m1[0]+m1[1];\n"
"float8 t2 = m2[0]+m2[1];\n"
"float8 t3 = m3[0]+m3[1];\n"
"float4 t14 = t1.lo + t1.hi;\n"
"float4 t24 = t2.lo + t2.hi;\n"
"float4 t34 = t3.lo + t3.hi;\n"
"smem1[0] = t14.x+t14.y+t14.z+t14.w;\n"
"smem2[0] = t24.x+t24.y+t24.z+t24.w;\n"
"smem3[0] = t34.x+t34.y+t34.z+t34.w;\n"
"}\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"}\n"
"inline void reduce2(float val1, float val2, __local volatile float* smem1, __local volatile float* smem2, int tid)\n"
"{\n"
"smem1[tid] = val1;\n"
"smem2[tid] = val2;\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"if (tid < 32)\n"
"{\n"
"smem1[tid] += smem1[tid + 32];\n"
"smem2[tid] += smem2[tid + 32];\n"
"#if WAVE_SIZE < 32\n"
"}\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"if (tid < 16)\n"
"{\n"
"#endif\n"
"smem1[tid] += smem1[tid + 16];\n"
"smem2[tid] += smem2[tid + 16];\n"
"#if WAVE_SIZE <16\n"
"}\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"if (tid<1)\n"
"{\n"
"#endif\n"
"local float8* m1 = (local float8*)smem1;\n"
"local float8* m2 = (local float8*)smem2;\n"
"float8 t1 = m1[0]+m1[1];\n"
"float8 t2 = m2[0]+m2[1];\n"
"float4 t14 = t1.lo + t1.hi;\n"
"float4 t24 = t2.lo + t2.hi;\n"
"smem1[0] = t14.x+t14.y+t14.z+t14.w;\n"
"smem2[0] = t24.x+t24.y+t24.z+t24.w;\n"
"}\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"}\n"
"inline void reduce1(float val1, __local volatile float* smem1, int tid)\n"
"{\n"
"smem1[tid] = val1;\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"if (tid < 32)\n"
"{\n"
"smem1[tid] += smem1[tid + 32];\n"
"#if WAVE_SIZE < 32\n"
"}\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"if (tid < 16)\n"
"{\n"
"#endif\n"
"smem1[tid] += smem1[tid + 16];\n"
"#if WAVE_SIZE <16\n"
"}\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"if (tid<1)\n"
"{\n"
"#endif\n"
"local float8* m1 = (local float8*)smem1;\n"
"float8 t1 = m1[0]+m1[1];\n"
"float4 t14 = t1.lo + t1.hi;\n"
"smem1[0] = t14.x+t14.y+t14.z+t14.w;\n"
"}\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"}\n"
"#endif\n"
"#define SCALE (1.0f / (1 << 20))\n"
"#define  THRESHOLD  0.01f\n"
"__constant sampler_t sampler    = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_LINEAR;\n"
"#define VAL(_y,_x,_yy,_xx)    (IPatchLocal[(yid+((_y)*LSy)+1+(_yy))*LM_W+(xid+((_x)*LSx)+1+(_xx))])\n"
"inline void SetPatch(local float* IPatchLocal, int TileY, int TileX,\n"
"float* Pch, float* Dx, float* Dy,\n"
"float* A11, float* A12, float* A22, float w)\n"
"{\n"
"unsigned int xid=get_local_id(0);\n"
"unsigned int yid=get_local_id(1);\n"
"*Pch = VAL(TileY,TileX,0,0);\n"
"float dIdx = (3.0f*VAL(TileY,TileX,-1,1)+10.0f*VAL(TileY,TileX,0,1)+3.0f*VAL(TileY,TileX,+1,1))-(3.0f*VAL(TileY,TileX,-1,-1)+10.0f*VAL(TileY,TileX,0,-1)+3.0f*VAL(TileY,TileX,+1,-1));\n"
"float dIdy = (3.0f*VAL(TileY,TileX,1,-1)+10.0f*VAL(TileY,TileX,1,0)+3.0f*VAL(TileY,TileX,1,+1))-(3.0f*VAL(TileY,TileX,-1,-1)+10.0f*VAL(TileY,TileX,-1,0)+3.0f*VAL(TileY,TileX,-1,+1));\n"
"dIdx *= w;\n"
"dIdy *= w;\n"
"*Dx = dIdx;\n"
"*Dy = dIdy;\n"
"*A11 += dIdx * dIdx;\n"
"*A12 += dIdx * dIdy;\n"
"*A22 += dIdy * dIdy;\n"
"}\n"
"#undef VAL\n"
"inline void GetPatch(image2d_t J, float x, float y,\n"
"float* Pch, float* Dx, float* Dy,\n"
"float* b1, float* b2)\n"
"{\n"
"float J_val = read_imagef(J, sampler, (float2)(x, y)).x;\n"
"float diff = (J_val - *Pch) * 32.0f;\n"
"*b1 += diff**Dx;\n"
"*b2 += diff**Dy;\n"
"}\n"
"inline void GetError(image2d_t J, const float x, const float y, const float* Pch, float* errval)\n"
"{\n"
"float diff = read_imagef(J, sampler, (float2)(x,y)).x-*Pch;\n"
"*errval += fabs(diff);\n"
"}\n"
"#define READI(_y,_x)    IPatchLocal[(yid+((_y)*LSy))*LM_W+(xid+((_x)*LSx))] = read_imagef(I, sampler, (float2)(Point.x + xid+(_x)*LSx + 0.5f-1, Point.y + yid+(_y)*LSy+ 0.5f-1)).x;\n"
"void ReadPatchIToLocalMem(image2d_t I, float2 Point, local float* IPatchLocal)\n"
"{\n"
"unsigned int xid=get_local_id(0);\n"
"unsigned int yid=get_local_id(1);\n"
"READI(0,0);READI(0,1);READI(0,2);\n"
"READI(1,0);READI(1,1);READI(1,2);\n"
"READI(2,0);READI(2,1);READI(2,2);\n"
"if(xid<2)\n"
"{\n"
"READI(0,3);\n"
"READI(1,3);\n"
"READI(2,3);\n"
"}\n"
"if(yid<2)\n"
"{\n"
"READI(3,0);READI(3,1);READI(3,2);\n"
"}\n"
"if(yid<2 && xid<2)\n"
"{\n"
"READI(3,3);\n"
"}\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"}\n"
"#undef READI\n"
"__attribute__((reqd_work_group_size(LSx, LSy, 1)))\n"
"__kernel void lkSparse(image2d_t I, image2d_t J,\n"
"__global const float2* prevPts, __global float2* nextPts, __global uchar* status, __global float* err,\n"
"const int level, const int rows, const int cols, int PATCH_X, int PATCH_Y, int c_winSize_x, int c_winSize_y, int c_iters, char calcErr)\n"
"{\n"
"__local float smem1[BUFFER];\n"
"__local float smem2[BUFFER];\n"
"__local float smem3[BUFFER];\n"
"unsigned int xid=get_local_id(0);\n"
"unsigned int yid=get_local_id(1);\n"
"unsigned int gid=get_group_id(0);\n"
"unsigned int xsize=get_local_size(0);\n"
"unsigned int ysize=get_local_size(1);\n"
"int xBase, yBase, k;\n"
"float wx = ((xid+2*xsize)<c_winSize_x)?1:0;\n"
"float wy = ((yid+2*ysize)<c_winSize_y)?1:0;\n"
"float2 c_halfWin = (float2)((c_winSize_x - 1)>>1, (c_winSize_y - 1)>>1);\n"
"const int tid = mad24(yid, xsize, xid);\n"
"float2 prevPt = prevPts[gid] / (float2)(1 << level);\n"
"if (prevPt.x < 0 || prevPt.x >= cols || prevPt.y < 0 || prevPt.y >= rows)\n"
"{\n"
"if (tid == 0 && level == 0)\n"
"{\n"
"status[gid] = 0;\n"
"}\n"
"return;\n"
"}\n"
"prevPt -= c_halfWin;\n"
"float A11 = 0;\n"
"float A12 = 0;\n"
"float A22 = 0;\n"
"float I_patch[GRIDSIZE][GRIDSIZE];\n"
"float dIdx_patch[GRIDSIZE][GRIDSIZE];\n"
"float dIdy_patch[GRIDSIZE][GRIDSIZE];\n"
"local float IPatchLocal[LM_W*LM_H];\n"
"ReadPatchIToLocalMem(I,prevPt,IPatchLocal);\n"
"{\n"
"SetPatch(IPatchLocal, 0, 0,\n"
"&I_patch[0][0], &dIdx_patch[0][0], &dIdy_patch[0][0],\n"
"&A11, &A12, &A22,1);\n"
"SetPatch(IPatchLocal, 0, 1,\n"
"&I_patch[0][1], &dIdx_patch[0][1], &dIdy_patch[0][1],\n"
"&A11, &A12, &A22,1);\n"
"SetPatch(IPatchLocal, 0, 2,\n"
"&I_patch[0][2], &dIdx_patch[0][2], &dIdy_patch[0][2],\n"
"&A11, &A12, &A22,wx);\n"
"}\n"
"{\n"
"SetPatch(IPatchLocal, 1, 0,\n"
"&I_patch[1][0], &dIdx_patch[1][0], &dIdy_patch[1][0],\n"
"&A11, &A12, &A22,1);\n"
"SetPatch(IPatchLocal, 1,1,\n"
"&I_patch[1][1], &dIdx_patch[1][1], &dIdy_patch[1][1],\n"
"&A11, &A12, &A22,1);\n"
"SetPatch(IPatchLocal, 1,2,\n"
"&I_patch[1][2], &dIdx_patch[1][2], &dIdy_patch[1][2],\n"
"&A11, &A12, &A22,wx);\n"
"}\n"
"{\n"
"SetPatch(IPatchLocal, 2,0,\n"
"&I_patch[2][0], &dIdx_patch[2][0], &dIdy_patch[2][0],\n"
"&A11, &A12, &A22,wy);\n"
"SetPatch(IPatchLocal, 2,1,\n"
"&I_patch[2][1], &dIdx_patch[2][1], &dIdy_patch[2][1],\n"
"&A11, &A12, &A22,wy);\n"
"SetPatch(IPatchLocal, 2,2,\n"
"&I_patch[2][2], &dIdx_patch[2][2], &dIdy_patch[2][2],\n"
"&A11, &A12, &A22,wx*wy);\n"
"}\n"
"reduce3(A11, A12, A22, smem1, smem2, smem3, tid);\n"
"A11 = smem1[0];\n"
"A12 = smem2[0];\n"
"A22 = smem3[0];\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"float D = A11 * A22 - A12 * A12;\n"
"if (D < 1.192092896e-07f)\n"
"{\n"
"if (tid == 0 && level == 0)\n"
"status[gid] = 0;\n"
"return;\n"
"}\n"
"A11 /= D;\n"
"A12 /= D;\n"
"A22 /= D;\n"
"prevPt = nextPts[gid] * 2.0f - c_halfWin;\n"
"for (k = 0; k < c_iters; ++k)\n"
"{\n"
"if (prevPt.x < -c_halfWin.x || prevPt.x >= cols || prevPt.y < -c_halfWin.y || prevPt.y >= rows)\n"
"{\n"
"if (tid == 0 && level == 0)\n"
"status[gid] = 0;\n"
"break;\n"
"}\n"
"float b1 = 0;\n"
"float b2 = 0;\n"
"yBase=yid;\n"
"{\n"
"xBase=xid;\n"
"GetPatch(J, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f,\n"
"&I_patch[0][0], &dIdx_patch[0][0], &dIdy_patch[0][0],\n"
"&b1, &b2);\n"
"xBase+=xsize;\n"
"GetPatch(J, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f,\n"
"&I_patch[0][1], &dIdx_patch[0][1], &dIdy_patch[0][1],\n"
"&b1, &b2);\n"
"xBase+=xsize;\n"
"GetPatch(J, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f,\n"
"&I_patch[0][2], &dIdx_patch[0][2], &dIdy_patch[0][2],\n"
"&b1, &b2);\n"
"}\n"
"yBase+=ysize;\n"
"{\n"
"xBase=xid;\n"
"GetPatch(J, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f,\n"
"&I_patch[1][0], &dIdx_patch[1][0], &dIdy_patch[1][0],\n"
"&b1, &b2);\n"
"xBase+=xsize;\n"
"GetPatch(J, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f,\n"
"&I_patch[1][1], &dIdx_patch[1][1], &dIdy_patch[1][1],\n"
"&b1, &b2);\n"
"xBase+=xsize;\n"
"GetPatch(J, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f,\n"
"&I_patch[1][2], &dIdx_patch[1][2], &dIdy_patch[1][2],\n"
"&b1, &b2);\n"
"}\n"
"yBase+=ysize;\n"
"{\n"
"xBase=xid;\n"
"GetPatch(J, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f,\n"
"&I_patch[2][0], &dIdx_patch[2][0], &dIdy_patch[2][0],\n"
"&b1, &b2);\n"
"xBase+=xsize;\n"
"GetPatch(J, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f,\n"
"&I_patch[2][1], &dIdx_patch[2][1], &dIdy_patch[2][1],\n"
"&b1, &b2);\n"
"xBase+=xsize;\n"
"GetPatch(J, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f,\n"
"&I_patch[2][2], &dIdx_patch[2][2], &dIdy_patch[2][2],\n"
"&b1, &b2);\n"
"}\n"
"reduce2(b1, b2, smem1, smem2, tid);\n"
"b1 = smem1[0];\n"
"b2 = smem2[0];\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"float2 delta;\n"
"delta.x = A12 * b2 - A22 * b1;\n"
"delta.y = A12 * b1 - A11 * b2;\n"
"prevPt += delta;\n"
"if (fabs(delta.x) < THRESHOLD && fabs(delta.y) < THRESHOLD)\n"
"break;\n"
"}\n"
"D = 0.0f;\n"
"if (calcErr)\n"
"{\n"
"yBase=yid;\n"
"{\n"
"xBase=xid;\n"
"GetError(J, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f,\n"
"&I_patch[0][0], &D);\n"
"xBase+=xsize;\n"
"GetError(J, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f,\n"
"&I_patch[0][1], &D);\n"
"xBase+=xsize;\n"
"if(xBase<c_winSize_x)\n"
"GetError(J, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f,\n"
"&I_patch[0][2], &D);\n"
"}\n"
"yBase+=ysize;\n"
"{\n"
"xBase=xid;\n"
"GetError(J, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f,\n"
"&I_patch[1][0], &D);\n"
"xBase+=xsize;\n"
"GetError(J, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f,\n"
"&I_patch[1][1], &D);\n"
"xBase+=xsize;\n"
"if(xBase<c_winSize_x)\n"
"GetError(J, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f,\n"
"&I_patch[1][2], &D);\n"
"}\n"
"yBase+=ysize;\n"
"if(yBase<c_winSize_y)\n"
"{\n"
"xBase=xid;\n"
"GetError(J, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f,\n"
"&I_patch[2][0], &D);\n"
"xBase+=xsize;\n"
"GetError(J, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f,\n"
"&I_patch[2][1], &D);\n"
"xBase+=xsize;\n"
"if(xBase<c_winSize_x)\n"
"GetError(J, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f,\n"
"&I_patch[2][2], &D);\n"
"}\n"
"reduce1(D, smem1, tid);\n"
"}\n"
"if (tid == 0)\n"
"{\n"
"prevPt += c_halfWin;\n"
"nextPts[gid] = prevPt;\n"
"if (calcErr)\n"
"err[gid] = smem1[0] / (float)(c_winSize_x * c_winSize_y);\n"
"}\n"
"}\n"
, "b7099fcbc60bd5528dacc491eadd88c1"};
ProgramSource pyrlk_oclsrc(pyrlk.programStr);
}
}}