| // 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); |
| } |
| }} |
