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

#include "opencv2/core.hpp"
#include "cvconfig.h"
#include "opencl_kernels_xfeatures2d.hpp"

#ifdef HAVE_OPENCL

namespace cv
{
namespace ocl
{
namespace xfeatures2d
{

static const char* const moduleName = "xfeatures2d";

struct cv::ocl::internal::ProgramEntry surf_oclsrc={moduleName, "surf",
"#define ORI_SEARCH_INC  5\n"
"#define ORI_LOCAL_SIZE  (360 / ORI_SEARCH_INC)\n"
"#ifndef HAVE_IMAGE2D\n"
"__inline uint read_sumTex_(__global uint* sumTex, int sum_step, int img_rows, int img_cols, int2 coord)\n"
"{\n"
"int x = clamp(coord.x, 0, img_cols);\n"
"int y = clamp(coord.y, 0, img_rows);\n"
"return sumTex[sum_step * y + x];\n"
"}\n"
"__inline uchar read_imgTex_(__global uchar* imgTex, int img_step, int img_rows, int img_cols, float2 coord)\n"
"{\n"
"int x = clamp(convert_int_rte(coord.x), 0, img_cols-1);\n"
"int y = clamp(convert_int_rte(coord.y), 0, img_rows-1);\n"
"return imgTex[img_step * y + x];\n"
"}\n"
"#define read_sumTex(coord) read_sumTex_(sumTex, sum_step, img_rows, img_cols, coord)\n"
"#define read_imgTex(coord) read_imgTex_(imgTex, img_step, img_rows, img_cols, coord)\n"
"#define __PARAM_sumTex__ __global uint* sumTex, int sum_step, int sum_offset\n"
"#define __PARAM_imgTex__ __global uchar* imgTex, int img_step, int img_offset\n"
"#define __PASS_sumTex__ sumTex, sum_step, sum_offset\n"
"#define __PASS_imgTex__ imgTex, img_step, img_offset\n"
"#else\n"
"__inline uint read_sumTex_(image2d_t sumTex, sampler_t sam, int2 coord)\n"
"{\n"
"return read_imageui(sumTex, sam, coord).x;\n"
"}\n"
"__inline uchar read_imgTex_(image2d_t imgTex, sampler_t sam, float2 coord)\n"
"{\n"
"return (uchar)read_imageui(imgTex, sam, coord).x;\n"
"}\n"
"#define read_sumTex(coord) read_sumTex_(sumTex, sampler, coord)\n"
"#define read_imgTex(coord) read_imgTex_(imgTex, sampler, coord)\n"
"#define __PARAM_sumTex__ image2d_t sumTex\n"
"#define __PARAM_imgTex__ image2d_t imgTex\n"
"#define __PASS_sumTex__ sumTex\n"
"#define __PASS_imgTex__ imgTex\n"
"#endif\n"
"#if defined (DOUBLE_SUPPORT)\n"
"#ifdef cl_khr_fp64\n"
"#pragma OPENCL EXTENSION cl_khr_fp64:enable\n"
"#elif defined (cl_amd_fp64)\n"
"#pragma OPENCL EXTENSION cl_amd_fp64:enable\n"
"#endif\n"
"#define F double\n"
"#else\n"
"#define F float\n"
"#endif\n"
"__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;\n"
"#ifndef FLT_EPSILON\n"
"#define FLT_EPSILON (1e-15)\n"
"#endif\n"
"#ifndef CV_PI_F\n"
"#define CV_PI_F 3.14159265f\n"
"#endif\n"
"__inline int calcSize(int octave, int layer)\n"
"{\n"
"const int HAAR_SIZE0 = 9;\n"
"const int HAAR_SIZE_INC = 6;\n"
"return (HAAR_SIZE0 + HAAR_SIZE_INC * layer) << octave;\n"
"}\n"
"F calcAxisAlignedDerivative(\n"
"int plus1a_A, int plus1a_B, int plus1a_C, int plus1a_D, F plus1a_scale,\n"
"int plus1b_A, int plus1b_B, int plus1b_C, int plus1b_D, F plus1b_scale,\n"
"int minus2_A, int minus2_B, int minus2_C, int minus2_D, F minus2_scale)\n"
"{\n"
"F plus1a = plus1a_A - plus1a_B - plus1a_C + plus1a_D;\n"
"F plus1b = plus1b_A - plus1b_B - plus1b_C + plus1b_D;\n"
"F minus2 = minus2_A - minus2_B - minus2_C + minus2_D;\n"
"return (plus1a / plus1a_scale -\n"
"2.0f * minus2 / minus2_scale +\n"
"plus1b / plus1b_scale);\n"
"}\n"
"__kernel void SURF_calcLayerDetAndTrace(\n"
"__PARAM_sumTex__,\n"
"int img_rows, int img_cols,\n"
"int c_nOctaveLayers, int c_octave, int c_layer_rows,\n"
"__global float * det,\n"
"int det_step, int det_offset,\n"
"__global float * trace,\n"
"int trace_step, int trace_offset)\n"
"{\n"
"det_step   /= sizeof(*det);\n"
"trace_step /= sizeof(*trace);\n"
"#ifndef HAVE_IMAGE2D\n"
"sum_step/= sizeof(uint);\n"
"#endif\n"
"const int gridDim_y  = get_num_groups(1) / (c_nOctaveLayers + 2);\n"
"const int blockIdx_y = get_group_id(1) % gridDim_y;\n"
"const int blockIdx_z = get_group_id(1) / gridDim_y;\n"
"const int j = get_local_id(0) + get_group_id(0) * get_local_size(0);\n"
"const int i = get_local_id(1) + blockIdx_y * get_local_size(1);\n"
"const int layer = blockIdx_z;\n"
"const int size = calcSize(c_octave, layer);\n"
"const int samples_i = 1 + ((img_rows - size) >> c_octave);\n"
"const int samples_j = 1 + ((img_cols - size) >> c_octave);\n"
"const int margin = (size >> 1) >> c_octave;\n"
"if (size <= img_rows && size <= img_cols && i < samples_i && j < samples_j)\n"
"{\n"
"int x = j << c_octave;\n"
"int y = i << c_octave;\n"
"float ratio = (float)size / 9;\n"
"int r1 = round(ratio);\n"
"int r2 = round(ratio * 2.0f);\n"
"int r3 = round(ratio * 3.0f);\n"
"int r4 = round(ratio * 4.0f);\n"
"int r5 = round(ratio * 5.0f);\n"
"int r6 = round(ratio * 6.0f);\n"
"int r7 = round(ratio * 7.0f);\n"
"int r8 = round(ratio * 8.0f);\n"
"int r9 = round(ratio * 9.0f);\n"
"F d = 0;\n"
"{\n"
"int t02 = read_sumTex( (int2)(x, y + r2));\n"
"int t07 = read_sumTex( (int2)(x, y + r7));\n"
"int t32 = read_sumTex( (int2)(x + r3, y + r2));\n"
"int t37 = read_sumTex( (int2)(x + r3, y + r7));\n"
"int t62 = read_sumTex( (int2)(x + r6, y + r2));\n"
"int t67 = read_sumTex( (int2)(x + r6, y + r7));\n"
"int t92 = read_sumTex( (int2)(x + r9, y + r2));\n"
"int t97 = read_sumTex( (int2)(x + r9, y + r7));\n"
"d = calcAxisAlignedDerivative(t02, t07, t32, t37, (r3) * (r7 - r2),\n"
"t62, t67, t92, t97, (r9 - r6) * (r7 - r2),\n"
"t32, t37, t62, t67, (r6 - r3) * (r7 - r2));\n"
"}\n"
"const float dx  = (float)d;\n"
"d = 0;\n"
"{\n"
"int t20 = read_sumTex( (int2)(x + r2, y) );\n"
"int t23 = read_sumTex( (int2)(x + r2, y + r3) );\n"
"int t70 = read_sumTex( (int2)(x + r7, y) );\n"
"int t73 = read_sumTex( (int2)(x + r7, y + r3) );\n"
"int t26 = read_sumTex( (int2)(x + r2, y + r6) );\n"
"int t76 = read_sumTex( (int2)(x + r7, y + r6) );\n"
"int t29 = read_sumTex( (int2)(x + r2, y + r9) );\n"
"int t79 = read_sumTex( (int2)(x + r7, y + r9) );\n"
"d = calcAxisAlignedDerivative(t20, t23, t70, t73, (r7 - r2) * (r3),\n"
"t26, t29, t76, t79, (r7 - r2) * (r9 - r6),\n"
"t23, t26, t73, t76, (r7 - r2) * (r6 - r3));\n"
"}\n"
"const float dy  = (float)d;\n"
"d = 0;\n"
"{\n"
"F t = 0;\n"
"t += read_sumTex( (int2)(x + r1, y + r1) );\n"
"t -= read_sumTex( (int2)(x + r1, y + r4) );\n"
"t -= read_sumTex( (int2)(x + r4, y + r1) );\n"
"t += read_sumTex( (int2)(x + r4, y + r4) );\n"
"d += t / ((r4 - r1) * (r4 - r1));\n"
"t = 0;\n"
"t += read_sumTex( (int2)(x + r5, y + r1) );\n"
"t -= read_sumTex( (int2)(x + r5, y + r4) );\n"
"t -= read_sumTex( (int2)(x + r8, y + r1) );\n"
"t += read_sumTex( (int2)(x + r8, y + r4) );\n"
"d -= t / ((r8 - r5) * (r4 - r1));\n"
"t = 0;\n"
"t += read_sumTex( (int2)(x + r1, y + r5) );\n"
"t -= read_sumTex( (int2)(x + r1, y + r8) );\n"
"t -= read_sumTex( (int2)(x + r4, y + r5) );\n"
"t += read_sumTex( (int2)(x + r4, y + r8) );\n"
"d -= t / ((r4 - r1) * (r8 - r5));\n"
"t = 0;\n"
"t += read_sumTex( (int2)(x + r5, y + r5) );\n"
"t -= read_sumTex( (int2)(x + r5, y + r8) );\n"
"t -= read_sumTex( (int2)(x + r8, y + r5) );\n"
"t += read_sumTex( (int2)(x + r8, y + r8) );\n"
"d += t / ((r8 - r5) * (r8 - r5));\n"
"}\n"
"const float dxy = (float)d;\n"
"det  [j + margin + det_step   * (layer * c_layer_rows + i + margin)] = dx * dy - 0.81f * dxy * dxy;\n"
"trace[j + margin + trace_step * (layer * c_layer_rows + i + margin)] = dx + dy;\n"
"}\n"
"}\n"
"__kernel\n"
"void SURF_findMaximaInLayer(\n"
"__global float * det,\n"
"int det_step, int det_offset,\n"
"__global float * trace,\n"
"int trace_step, int trace_offset,\n"
"__global int4 * maxPosBuffer,\n"
"volatile __global  int* maxCounter,\n"
"int counter_offset,\n"
"int img_rows,\n"
"int img_cols,\n"
"int c_nOctaveLayers,\n"
"int c_octave,\n"
"int c_layer_rows,\n"
"int c_layer_cols,\n"
"int c_max_candidates,\n"
"float c_hessianThreshold\n"
")\n"
"{\n"
"volatile __local  float N9[768];\n"
"det_step   /= sizeof(float);\n"
"trace_step /= sizeof(float);\n"
"maxCounter += counter_offset;\n"
"const int gridDim_y  = get_num_groups(1) / c_nOctaveLayers;\n"
"const int blockIdx_y = get_group_id(1)   % gridDim_y;\n"
"const int blockIdx_z = get_group_id(1)   / gridDim_y;\n"
"const int layer = blockIdx_z + 1;\n"
"const int size = calcSize(c_octave, layer);\n"
"const int margin = ((calcSize(c_octave, layer + 1) >> 1) >> c_octave) + 1;\n"
"const int j = get_local_id(0) + get_group_id(0) * (get_local_size(0) - 2) + margin - 1;\n"
"const int i = get_local_id(1) + blockIdx_y      * (get_local_size(1) - 2) + margin - 1;\n"
"const int zoff     = get_local_size(0) * get_local_size(1);\n"
"const int localLin = get_local_id(0) + get_local_id(1) * get_local_size(0) + zoff;\n"
"int l_x = min(max(j, 0), img_cols - 1);\n"
"int l_y = c_layer_rows * layer + min(max(i, 0), img_rows - 1);\n"
"N9[localLin - zoff] =\n"
"det[det_step * (l_y - c_layer_rows) + l_x];\n"
"N9[localLin       ] =\n"
"det[det_step * (l_y               ) + l_x];\n"
"N9[localLin + zoff] =\n"
"det[det_step * (l_y + c_layer_rows) + l_x];\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"if (i < c_layer_rows - margin\n"
"&& j < c_layer_cols - margin\n"
"&& get_local_id(0) > 0\n"
"&& get_local_id(0) < get_local_size(0) - 1\n"
"&& get_local_id(1) > 0\n"
"&& get_local_id(1) < get_local_size(1) - 1\n"
")\n"
"{\n"
"float val0 = N9[localLin];\n"
"if (val0 > c_hessianThreshold)\n"
"{\n"
"const bool condmax = val0 > N9[localLin - 1 - get_local_size(0) - zoff]\n"
"&&                   val0 > N9[localLin     - get_local_size(0) - zoff]\n"
"&&                   val0 > N9[localLin + 1 - get_local_size(0) - zoff]\n"
"&&                   val0 > N9[localLin - 1                     - zoff]\n"
"&&                   val0 > N9[localLin                         - zoff]\n"
"&&                   val0 > N9[localLin + 1                     - zoff]\n"
"&&                   val0 > N9[localLin - 1 + get_local_size(0) - zoff]\n"
"&&                   val0 > N9[localLin     + get_local_size(0) - zoff]\n"
"&&                   val0 > N9[localLin + 1 + get_local_size(0) - zoff]\n"
"&&                   val0 > N9[localLin - 1 - get_local_size(0)]\n"
"&&                   val0 > N9[localLin     - get_local_size(0)]\n"
"&&                   val0 > N9[localLin + 1 - get_local_size(0)]\n"
"&&                   val0 > N9[localLin - 1                    ]\n"
"&&                   val0 > N9[localLin + 1                    ]\n"
"&&                   val0 > N9[localLin - 1 + get_local_size(0)]\n"
"&&                   val0 > N9[localLin     + get_local_size(0)]\n"
"&&                   val0 > N9[localLin + 1 + get_local_size(0)]\n"
"&&                   val0 > N9[localLin - 1 - get_local_size(0) + zoff]\n"
"&&                   val0 > N9[localLin     - get_local_size(0) + zoff]\n"
"&&                   val0 > N9[localLin + 1 - get_local_size(0) + zoff]\n"
"&&                   val0 > N9[localLin - 1                     + zoff]\n"
"&&                   val0 > N9[localLin                         + zoff]\n"
"&&                   val0 > N9[localLin + 1                     + zoff]\n"
"&&                   val0 > N9[localLin - 1 + get_local_size(0) + zoff]\n"
"&&                   val0 > N9[localLin     + get_local_size(0) + zoff]\n"
"&&                   val0 > N9[localLin + 1 + get_local_size(0) + zoff]\n"
";\n"
"if(condmax)\n"
"{\n"
"int ind = atomic_inc(maxCounter);\n"
"if (ind < c_max_candidates)\n"
"{\n"
"const int laplacian = (int) copysign(1.0f, trace[trace_step* (layer * c_layer_rows + i) + j]);\n"
"maxPosBuffer[ind] = (int4)(j, i, layer, laplacian);\n"
"}\n"
"}\n"
"}\n"
"}\n"
"}\n"
"inline bool solve3x3_float(const float4 *A, const float *b, float *x)\n"
"{\n"
"float det = A[0].x * (A[1].y * A[2].z - A[1].z * A[2].y)\n"
"- A[0].y * (A[1].x * A[2].z - A[1].z * A[2].x)\n"
"+ A[0].z * (A[1].x * A[2].y - A[1].y * A[2].x);\n"
"if (det != 0)\n"
"{\n"
"F invdet = 1.0f / det;\n"
"x[0] = invdet *\n"
"(b[0]    * (A[1].y * A[2].z - A[1].z * A[2].y) -\n"
"A[0].y * (b[1]    * A[2].z - A[1].z * b[2]   ) +\n"
"A[0].z * (b[1]    * A[2].y - A[1].y * b[2]   ));\n"
"x[1] = invdet *\n"
"(A[0].x * (b[1]    * A[2].z - A[1].z * b[2]   ) -\n"
"b[0]    * (A[1].x * A[2].z - A[1].z * A[2].x) +\n"
"A[0].z * (A[1].x * b[2]    - b[1]    * A[2].x));\n"
"x[2] = invdet *\n"
"(A[0].x * (A[1].y * b[2]    - b[1]    * A[2].y) -\n"
"A[0].y * (A[1].x * b[2]    - b[1]    * A[2].x) +\n"
"b[0]    * (A[1].x * A[2].y - A[1].y * A[2].x));\n"
"return true;\n"
"}\n"
"return false;\n"
"}\n"
"#define X_ROW          0\n"
"#define Y_ROW          1\n"
"#define LAPLACIAN_ROW  2\n"
"#define OCTAVE_ROW     3\n"
"#define SIZE_ROW       4\n"
"#define ANGLE_ROW      5\n"
"#define HESSIAN_ROW    6\n"
"#define ROWS_COUNT     7\n"
"__kernel\n"
"void SURF_interpolateKeypoint(\n"
"__global const float * det,\n"
"int det_step, int det_offset,\n"
"__global const int4 * maxPosBuffer,\n"
"__global float * keypoints,\n"
"int keypoints_step, int keypoints_offset,\n"
"volatile __global int* featureCounter,\n"
"int img_rows,\n"
"int img_cols,\n"
"int c_octave,\n"
"int c_layer_rows,\n"
"int c_max_features\n"
")\n"
"{\n"
"det_step /= sizeof(*det);\n"
"keypoints_step /= sizeof(*keypoints);\n"
"__global float * featureX       = keypoints + X_ROW * keypoints_step;\n"
"__global float * featureY       = keypoints + Y_ROW * keypoints_step;\n"
"__global int * featureLaplacian = (__global int *)keypoints + LAPLACIAN_ROW * keypoints_step;\n"
"__global int * featureOctave    = (__global int *)keypoints + OCTAVE_ROW * keypoints_step;\n"
"__global float * featureSize    = keypoints + SIZE_ROW * keypoints_step;\n"
"__global float * featureHessian = keypoints + HESSIAN_ROW * keypoints_step;\n"
"const int4 maxPos = maxPosBuffer[get_group_id(0)];\n"
"const int j = maxPos.x - 1 + get_local_id(0);\n"
"const int i = maxPos.y - 1 + get_local_id(1);\n"
"const int layer = maxPos.z - 1 + get_local_id(2);\n"
"volatile __local  float N9[3][3][3];\n"
"N9[get_local_id(2)][get_local_id(1)][get_local_id(0)] =\n"
"det[det_step * (c_layer_rows * layer + i) + j];\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"if (get_local_id(0) == 0 && get_local_id(1) == 0 && get_local_id(2) == 0)\n"
"{\n"
"float dD[3];\n"
"dD[0] = -0.5f * (N9[1][1][2] - N9[1][1][0]);\n"
"dD[1] = -0.5f * (N9[1][2][1] - N9[1][0][1]);\n"
"dD[2] = -0.5f * (N9[2][1][1] - N9[0][1][1]);\n"
"float4 H[3];\n"
"H[0].x = N9[1][1][0] - 2.0f * N9[1][1][1] + N9[1][1][2];\n"
"H[0].y= 0.25f * (N9[1][2][2] - N9[1][2][0] - N9[1][0][2] + N9[1][0][0]);\n"
"H[0].z= 0.25f * (N9[2][1][2] - N9[2][1][0] - N9[0][1][2] + N9[0][1][0]);\n"
"H[1].x = H[0].y;\n"
"H[1].y = N9[1][0][1] - 2.0f * N9[1][1][1] + N9[1][2][1];\n"
"H[1].z= 0.25f * (N9[2][2][1] - N9[2][0][1] - N9[0][2][1] + N9[0][0][1]);\n"
"H[2].x = H[0].z;\n"
"H[2].y = H[1].z;\n"
"H[2].z = N9[0][1][1] - 2.0f * N9[1][1][1] + N9[2][1][1];\n"
"float x[3];\n"
"if (solve3x3_float(H, dD, x))\n"
"{\n"
"if (fabs(x[0]) <= 1.f && fabs(x[1]) <= 1.f && fabs(x[2]) <= 1.f)\n"
"{\n"
"const int size = calcSize(c_octave, maxPos.z);\n"
"const int sum_i = (maxPos.y - ((size >> 1) >> c_octave)) << c_octave;\n"
"const int sum_j = (maxPos.x - ((size >> 1) >> c_octave)) << c_octave;\n"
"const float center_i = sum_i + (float)(size - 1) / 2;\n"
"const float center_j = sum_j + (float)(size - 1) / 2;\n"
"const float px = center_j + x[0] * (1 << c_octave);\n"
"const float py = center_i + x[1] * (1 << c_octave);\n"
"const int ds = size - calcSize(c_octave, maxPos.z - 1);\n"
"const float psize = round(size + x[2] * ds);\n"
"const float s = psize * 1.2f / 9.0f;\n"
"const int grad_wav_size = 2 * round(2.0f * s);\n"
"if ((img_rows + 1) >= grad_wav_size && (img_cols + 1) >= grad_wav_size)\n"
"{\n"
"int ind = atomic_inc(featureCounter);\n"
"if (ind < c_max_features)\n"
"{\n"
"featureX[ind] = px;\n"
"featureY[ind] = py;\n"
"featureLaplacian[ind] = maxPos.w;\n"
"featureOctave[ind] = c_octave;\n"
"featureSize[ind] = psize;\n"
"featureHessian[ind] = N9[1][1][1];\n"
"}\n"
"}\n"
"}\n"
"}\n"
"}\n"
"}\n"
"#define ORI_WIN       60\n"
"#define ORI_SAMPLES     113\n"
"#define ORI_RESPONSE_REDUCTION_WIDTH     48\n"
"#define ORI_RESPONSE_ARRAY_SIZE           (ORI_RESPONSE_REDUCTION_WIDTH * 2)\n"
"__constant float c_aptX[ORI_SAMPLES] = {-6, -5, -5, -5, -5, -5, -5, -5, -4, -4, -4, -4, -4, -4, -4, -4, -4, -3, -3, -3, -3, -3, -3, -3, -3, -3, -3, -3, -2, -2, -2, -2, -2, -2, -2, -2, -2, -2, -2, -1, -1, -1, -1, -1, -1, -1, -1, -1, -1, -1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 2, 2, 2, 2, 2, 2, 2, 2, 2, 2, 2, 3, 3, 3, 3, 3, 3, 3, 3, 3, 3, 3, 4, 4, 4, 4, 4, 4, 4, 4, 4, 5, 5, 5, 5, 5, 5, 5, 6};\n"
"__constant float c_aptY[ORI_SAMPLES] = {0, -3, -2, -1, 0, 1, 2, 3, -4, -3, -2, -1, 0, 1, 2, 3, 4, -5, -4, -3, -2, -1, 0, 1, 2, 3, 4, 5, -5, -4, -3, -2, -1, 0, 1, 2, 3, 4, 5, -5, -4, -3, -2, -1, 0, 1, 2, 3, 4, 5, -6, -5, -4, -3, -2, -1, 0, 1, 2, 3, 4, 5, 6, -5, -4, -3, -2, -1, 0, 1, 2, 3, 4, 5, -5, -4, -3, -2, -1, 0, 1, 2, 3, 4, 5, -5, -4, -3, -2, -1, 0, 1, 2, 3, 4, 5, -4, -3, -2, -1, 0, 1, 2, 3, 4, -3, -2, -1, 0, 1, 2, 3, 0};\n"
"__constant float c_aptW[ORI_SAMPLES] = {0.001455130288377404f, 0.001707611023448408f, 0.002547456417232752f, 0.003238451667129993f, 0.0035081731621176f,\n"
"0.003238451667129993f, 0.002547456417232752f, 0.001707611023448408f, 0.002003900473937392f, 0.0035081731621176f, 0.005233579315245152f,\n"
"0.00665318313986063f, 0.00720730796456337f, 0.00665318313986063f, 0.005233579315245152f, 0.0035081731621176f,\n"
"0.002003900473937392f, 0.001707611023448408f, 0.0035081731621176f, 0.006141661666333675f, 0.009162282571196556f,\n"
"0.01164754293859005f, 0.01261763460934162f, 0.01164754293859005f, 0.009162282571196556f, 0.006141661666333675f,\n"
"0.0035081731621176f, 0.001707611023448408f, 0.002547456417232752f, 0.005233579315245152f, 0.009162282571196556f,\n"
"0.01366852037608624f, 0.01737609319388866f, 0.0188232995569706f, 0.01737609319388866f, 0.01366852037608624f,\n"
"0.009162282571196556f, 0.005233579315245152f, 0.002547456417232752f, 0.003238451667129993f, 0.00665318313986063f,\n"
"0.01164754293859005f, 0.01737609319388866f, 0.02208934165537357f, 0.02392910048365593f, 0.02208934165537357f,\n"
"0.01737609319388866f, 0.01164754293859005f, 0.00665318313986063f, 0.003238451667129993f, 0.001455130288377404f,\n"
"0.0035081731621176f, 0.00720730796456337f, 0.01261763460934162f, 0.0188232995569706f, 0.02392910048365593f,\n"
"0.02592208795249462f, 0.02392910048365593f, 0.0188232995569706f, 0.01261763460934162f, 0.00720730796456337f,\n"
"0.0035081731621176f, 0.001455130288377404f, 0.003238451667129993f, 0.00665318313986063f, 0.01164754293859005f,\n"
"0.01737609319388866f, 0.02208934165537357f, 0.02392910048365593f, 0.02208934165537357f, 0.01737609319388866f,\n"
"0.01164754293859005f, 0.00665318313986063f, 0.003238451667129993f, 0.002547456417232752f, 0.005233579315245152f,\n"
"0.009162282571196556f, 0.01366852037608624f, 0.01737609319388866f, 0.0188232995569706f, 0.01737609319388866f,\n"
"0.01366852037608624f, 0.009162282571196556f, 0.005233579315245152f, 0.002547456417232752f, 0.001707611023448408f,\n"
"0.0035081731621176f, 0.006141661666333675f, 0.009162282571196556f, 0.01164754293859005f, 0.01261763460934162f,\n"
"0.01164754293859005f, 0.009162282571196556f, 0.006141661666333675f, 0.0035081731621176f, 0.001707611023448408f,\n"
"0.002003900473937392f, 0.0035081731621176f, 0.005233579315245152f, 0.00665318313986063f, 0.00720730796456337f,\n"
"0.00665318313986063f, 0.005233579315245152f, 0.0035081731621176f, 0.002003900473937392f, 0.001707611023448408f,\n"
"0.002547456417232752f, 0.003238451667129993f, 0.0035081731621176f, 0.003238451667129993f, 0.002547456417232752f,\n"
"0.001707611023448408f, 0.001455130288377404f\n"
"};\n"
"__constant float2 c_NX[5] = { (float2)(0, 2), (float2)(0, 0), (float2)(2, 4), (float2)(4, 4), (float2)(-1, 1) };\n"
"__constant float2 c_NY[5] = { (float2)(0, 0), (float2)(0, 2), (float2)(4, 4), (float2)(2, 4), (float2)(1, -1) };\n"
"void reduce_32_sum(volatile __local  float * data, volatile float* partial_reduction, int tid)\n"
"{\n"
"#define op(A, B) (*A)+(B)\n"
"data[tid] = *partial_reduction;\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"#ifndef WAVE_SIZE\n"
"#define WAVE_SIZE 1\n"
"#endif\n"
"if (tid < 16)\n"
"{\n"
"data[tid] = *partial_reduction = op(partial_reduction, data[tid + 16]);\n"
"#if WAVE_SIZE < 16\n"
"}\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"if (tid < 8)\n"
"{\n"
"#endif\n"
"data[tid] = *partial_reduction = op(partial_reduction, data[tid + 8]);\n"
"#if WAVE_SIZE < 8\n"
"}\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"if (tid < 4)\n"
"{\n"
"#endif\n"
"data[tid] = *partial_reduction = op(partial_reduction, data[tid + 4]);\n"
"#if WAVE_SIZE < 4\n"
"}\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"if (tid < 2)\n"
"{\n"
"#endif\n"
"data[tid] = *partial_reduction = op(partial_reduction, data[tid + 2 ]);\n"
"#if WAVE_SIZE < 2\n"
"}\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"if (tid < 1)\n"
"{\n"
"#endif\n"
"data[tid] = *partial_reduction = op(partial_reduction, data[tid + 1 ]);\n"
"}\n"
"#undef WAVE_SIZE\n"
"#undef op\n"
"}\n"
"__kernel\n"
"void SURF_calcOrientation(\n"
"__PARAM_sumTex__, int img_rows, int img_cols,\n"
"__global float * keypoints, int keypoints_step, int keypoints_offset )\n"
"{\n"
"keypoints_step /= sizeof(*keypoints);\n"
"#ifndef HAVE_IMAGE2D\n"
"sum_step       /= sizeof(uint);\n"
"#endif\n"
"__global float* featureX    = keypoints + X_ROW * keypoints_step;\n"
"__global float* featureY    = keypoints + Y_ROW * keypoints_step;\n"
"__global float* featureSize = keypoints + SIZE_ROW * keypoints_step;\n"
"__global float* featureDir  = keypoints + ANGLE_ROW * keypoints_step;\n"
"__local  float s_X[ORI_SAMPLES];\n"
"__local  float s_Y[ORI_SAMPLES];\n"
"__local  float s_angle[ORI_SAMPLES];\n"
"__local  float s_sumx[ORI_RESPONSE_ARRAY_SIZE];\n"
"__local  float s_sumy[ORI_RESPONSE_ARRAY_SIZE];\n"
"__local  float s_mod[ORI_RESPONSE_ARRAY_SIZE];\n"
"const float s = featureSize[get_group_id(0)] * 1.2f / 9.0f;\n"
"const int grad_wav_size = 2 * round(2.0f * s);\n"
"if ((img_rows + 1) < grad_wav_size || (img_cols + 1) < grad_wav_size)\n"
"return;\n"
"const int tid = get_local_id(0);\n"
"if (tid < ORI_RESPONSE_ARRAY_SIZE - ORI_LOCAL_SIZE) {\n"
"s_mod[tid + ORI_LOCAL_SIZE] = 0.0f;\n"
"}\n"
"float ratio = (float)grad_wav_size / 4;\n"
"int r2 = round(ratio * 2.0f);\n"
"int r4 = round(ratio * 4.0f);\n"
"for (int i = tid; i < ORI_SAMPLES; i += ORI_LOCAL_SIZE )\n"
"{\n"
"float X = 0.0f, Y = 0.0f, angle = 0.0f;\n"
"const float margin = (float)(grad_wav_size - 1) / 2.0f;\n"
"const int x = round(featureX[get_group_id(0)] + c_aptX[i] * s - margin);\n"
"const int y = round(featureY[get_group_id(0)] + c_aptY[i] * s - margin);\n"
"if (y >= 0 && y < (img_rows + 1) - grad_wav_size &&\n"
"x >= 0 && x < (img_cols + 1) - grad_wav_size)\n"
"{\n"
"float apt = c_aptW[i];\n"
"float t00 = read_sumTex( (int2)(x, y));\n"
"float t02 = read_sumTex( (int2)(x, y + r2));\n"
"float t04 = read_sumTex( (int2)(x, y + r4));\n"
"float t20 = read_sumTex( (int2)(x + r2, y));\n"
"float t24 = read_sumTex( (int2)(x + r2, y + r4));\n"
"float t40 = read_sumTex( (int2)(x + r4, y));\n"
"float t42 = read_sumTex( (int2)(x + r4, y + r2));\n"
"float t44 = read_sumTex( (int2)(x + r4, y + r4));\n"
"F t = t00 - t04 - t20 + t24;\n"
"X -= t / ((r2) * (r4));\n"
"t = t20 - t24 - t40 + t44;\n"
"X += t / ((r4 - r2) * (r4));\n"
"t = t00 - t02 - t40 + t42;\n"
"Y += t / ((r2) * (r4));\n"
"t = t02 - t04 - t42 + t44;\n"
"Y -= t  / ((r4) * (r4 - r2));\n"
"X = apt*X;\n"
"Y = apt*Y;\n"
"angle = atan2(Y, X);\n"
"if (angle < 0)\n"
"angle += 2.0f * CV_PI_F;\n"
"angle *= 180.0f / CV_PI_F;\n"
"}\n"
"s_X[i] = X;\n"
"s_Y[i] = Y;\n"
"s_angle[i] = angle;\n"
"}\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"float bestx = 0, besty = 0, best_mod = 0;\n"
"float sumx = 0.0f, sumy = 0.0f;\n"
"const int dir = tid * ORI_SEARCH_INC;\n"
"#pragma unroll\n"
"for (int i = 0; i < ORI_SAMPLES; ++i) {\n"
"int angle = round(s_angle[i]);\n"
"int d = abs(angle - dir);\n"
"if (d < ORI_WIN / 2 || d > 360 - ORI_WIN / 2)\n"
"{\n"
"sumx += s_X[i];\n"
"sumy += s_Y[i];\n"
"}\n"
"}\n"
"s_sumx[tid] = sumx;\n"
"s_sumy[tid] = sumy;\n"
"s_mod[tid] = sumx*sumx + sumy*sumy;\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"for(int t = ORI_RESPONSE_REDUCTION_WIDTH; t >= 3; t /= 2) {\n"
"if (tid < t) {\n"
"if (s_mod[tid] < s_mod[tid + t]) {\n"
"s_mod[tid] = s_mod[tid + t];\n"
"s_sumx[tid] = s_sumx[tid + t];\n"
"s_sumy[tid] = s_sumy[tid + t];\n"
"}\n"
"}\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"}\n"
"if (tid == 0)\n"
"{\n"
"int bestIdx = 0;\n"
"if (s_mod[1] > s_mod[bestIdx])\n"
"bestIdx = 1;\n"
"if (s_mod[2] > s_mod[bestIdx])\n"
"bestIdx = 2;\n"
"float kp_dir = atan2(s_sumy[bestIdx], s_sumx[bestIdx]);\n"
"if (kp_dir < 0)\n"
"kp_dir += 2.0f * CV_PI_F;\n"
"kp_dir *= 180.0f / CV_PI_F;\n"
"kp_dir = 360.0f - kp_dir;\n"
"if (fabs(kp_dir - 360.f) < FLT_EPSILON)\n"
"kp_dir = 0.f;\n"
"featureDir[get_group_id(0)] = kp_dir;\n"
"}\n"
"}\n"
"__kernel\n"
"void SURF_setUpRight(\n"
"__global float * keypoints,\n"
"int keypoints_step, int keypoints_offset,\n"
"int rows, int cols )\n"
"{\n"
"int i = get_global_id(0);\n"
"keypoints_step /= sizeof(*keypoints);\n"
"if(i < cols)\n"
"{\n"
"keypoints[mad24(keypoints_step, ANGLE_ROW, i)] = 270.f;\n"
"}\n"
"}\n"
"#undef ORI_SEARCH_INC\n"
"#undef ORI_WIN\n"
"#undef ORI_SAMPLES\n"
"#define PATCH_SZ 20\n"
"__constant float c_DW[PATCH_SZ * PATCH_SZ] =\n"
"{\n"
"3.695352233989979e-006f, 8.444558261544444e-006f, 1.760426494001877e-005f, 3.34794785885606e-005f, 5.808438800158911e-005f, 9.193058212986216e-005f, 0.0001327334757661447f, 0.0001748319627949968f, 0.0002100782439811155f, 0.0002302826324012131f, 0.0002302826324012131f, 0.0002100782439811155f, 0.0001748319627949968f, 0.0001327334757661447f, 9.193058212986216e-005f, 5.808438800158911e-005f, 3.34794785885606e-005f, 1.760426494001877e-005f, 8.444558261544444e-006f, 3.695352233989979e-006f,\n"
"8.444558261544444e-006f, 1.929736572492402e-005f, 4.022897701361217e-005f, 7.650675252079964e-005f, 0.0001327334903180599f, 0.0002100782585330308f, 0.0003033203829545528f, 0.0003995231236331165f, 0.0004800673632416874f, 0.0005262381164357066f, 0.0005262381164357066f, 0.0004800673632416874f, 0.0003995231236331165f, 0.0003033203829545528f, 0.0002100782585330308f, 0.0001327334903180599f, 7.650675252079964e-005f, 4.022897701361217e-005f, 1.929736572492402e-005f, 8.444558261544444e-006f,\n"
"1.760426494001877e-005f, 4.022897701361217e-005f, 8.386484114453197e-005f, 0.0001594926579855382f, 0.0002767078403849155f, 0.0004379475140012801f, 0.0006323281559161842f, 0.0008328808471560478f, 0.001000790391117334f, 0.001097041997127235f, 0.001097041997127235f, 0.001000790391117334f, 0.0008328808471560478f, 0.0006323281559161842f, 0.0004379475140012801f, 0.0002767078403849155f, 0.0001594926579855382f, 8.386484114453197e-005f, 4.022897701361217e-005f, 1.760426494001877e-005f,\n"
"3.34794785885606e-005f, 7.650675252079964e-005f, 0.0001594926579855382f, 0.0003033203247468919f, 0.0005262380582280457f, 0.0008328807889483869f, 0.001202550483867526f, 0.001583957928232849f, 0.001903285388834775f, 0.002086334861814976f, 0.002086334861814976f, 0.001903285388834775f, 0.001583957928232849f, 0.001202550483867526f, 0.0008328807889483869f, 0.0005262380582280457f, 0.0003033203247468919f, 0.0001594926579855382f, 7.650675252079964e-005f, 3.34794785885606e-005f,\n"
"5.808438800158911e-005f, 0.0001327334903180599f, 0.0002767078403849155f, 0.0005262380582280457f, 0.0009129836107604206f, 0.001444985857233405f, 0.002086335094645619f, 0.002748048631474376f, 0.00330205773934722f, 0.003619635012000799f, 0.003619635012000799f, 0.00330205773934722f, 0.002748048631474376f, 0.002086335094645619f, 0.001444985857233405f, 0.0009129836107604206f, 0.0005262380582280457f, 0.0002767078403849155f, 0.0001327334903180599f, 5.808438800158911e-005f,\n"
"9.193058212986216e-005f, 0.0002100782585330308f, 0.0004379475140012801f, 0.0008328807889483869f, 0.001444985857233405f, 0.002286989474669099f, 0.00330205773934722f, 0.004349356517195702f, 0.00522619066759944f, 0.005728822201490402f, 0.005728822201490402f, 0.00522619066759944f, 0.004349356517195702f, 0.00330205773934722f, 0.002286989474669099f, 0.001444985857233405f, 0.0008328807889483869f, 0.0004379475140012801f, 0.0002100782585330308f, 9.193058212986216e-005f,\n"
"0.0001327334757661447f, 0.0003033203829545528f, 0.0006323281559161842f, 0.001202550483867526f, 0.002086335094645619f, 0.00330205773934722f, 0.004767658654600382f, 0.006279794964939356f, 0.007545807864516974f, 0.008271530270576477f, 0.008271530270576477f, 0.007545807864516974f, 0.006279794964939356f, 0.004767658654600382f, 0.00330205773934722f, 0.002086335094645619f, 0.001202550483867526f, 0.0006323281559161842f, 0.0003033203829545528f, 0.0001327334757661447f,\n"
"0.0001748319627949968f, 0.0003995231236331165f, 0.0008328808471560478f, 0.001583957928232849f, 0.002748048631474376f, 0.004349356517195702f, 0.006279794964939356f, 0.008271529339253902f, 0.009939077310264111f, 0.01089497376233339f, 0.01089497376233339f, 0.009939077310264111f, 0.008271529339253902f, 0.006279794964939356f, 0.004349356517195702f, 0.002748048631474376f, 0.001583957928232849f, 0.0008328808471560478f, 0.0003995231236331165f, 0.0001748319627949968f,\n"
"0.0002100782439811155f, 0.0004800673632416874f, 0.001000790391117334f, 0.001903285388834775f, 0.00330205773934722f, 0.00522619066759944f, 0.007545807864516974f, 0.009939077310264111f, 0.01194280479103327f, 0.01309141051024199f, 0.01309141051024199f, 0.01194280479103327f, 0.009939077310264111f, 0.007545807864516974f, 0.00522619066759944f, 0.00330205773934722f, 0.001903285388834775f, 0.001000790391117334f, 0.0004800673632416874f, 0.0002100782439811155f,\n"
"0.0002302826324012131f, 0.0005262381164357066f, 0.001097041997127235f, 0.002086334861814976f, 0.003619635012000799f, 0.005728822201490402f, 0.008271530270576477f, 0.01089497376233339f, 0.01309141051024199f, 0.01435048412531614f, 0.01435048412531614f, 0.01309141051024199f, 0.01089497376233339f, 0.008271530270576477f, 0.005728822201490402f, 0.003619635012000799f, 0.002086334861814976f, 0.001097041997127235f, 0.0005262381164357066f, 0.0002302826324012131f,\n"
"0.0002302826324012131f, 0.0005262381164357066f, 0.001097041997127235f, 0.002086334861814976f, 0.003619635012000799f, 0.005728822201490402f, 0.008271530270576477f, 0.01089497376233339f, 0.01309141051024199f, 0.01435048412531614f, 0.01435048412531614f, 0.01309141051024199f, 0.01089497376233339f, 0.008271530270576477f, 0.005728822201490402f, 0.003619635012000799f, 0.002086334861814976f, 0.001097041997127235f, 0.0005262381164357066f, 0.0002302826324012131f,\n"
"0.0002100782439811155f, 0.0004800673632416874f, 0.001000790391117334f, 0.001903285388834775f, 0.00330205773934722f, 0.00522619066759944f, 0.007545807864516974f, 0.009939077310264111f, 0.01194280479103327f, 0.01309141051024199f, 0.01309141051024199f, 0.01194280479103327f, 0.009939077310264111f, 0.007545807864516974f, 0.00522619066759944f, 0.00330205773934722f, 0.001903285388834775f, 0.001000790391117334f, 0.0004800673632416874f, 0.0002100782439811155f,\n"
"0.0001748319627949968f, 0.0003995231236331165f, 0.0008328808471560478f, 0.001583957928232849f, 0.002748048631474376f, 0.004349356517195702f, 0.006279794964939356f, 0.008271529339253902f, 0.009939077310264111f, 0.01089497376233339f, 0.01089497376233339f, 0.009939077310264111f, 0.008271529339253902f, 0.006279794964939356f, 0.004349356517195702f, 0.002748048631474376f, 0.001583957928232849f, 0.0008328808471560478f, 0.0003995231236331165f, 0.0001748319627949968f,\n"
"0.0001327334757661447f, 0.0003033203829545528f, 0.0006323281559161842f, 0.001202550483867526f, 0.002086335094645619f, 0.00330205773934722f, 0.004767658654600382f, 0.006279794964939356f, 0.007545807864516974f, 0.008271530270576477f, 0.008271530270576477f, 0.007545807864516974f, 0.006279794964939356f, 0.004767658654600382f, 0.00330205773934722f, 0.002086335094645619f, 0.001202550483867526f, 0.0006323281559161842f, 0.0003033203829545528f, 0.0001327334757661447f,\n"
"9.193058212986216e-005f, 0.0002100782585330308f, 0.0004379475140012801f, 0.0008328807889483869f, 0.001444985857233405f, 0.002286989474669099f, 0.00330205773934722f, 0.004349356517195702f, 0.00522619066759944f, 0.005728822201490402f, 0.005728822201490402f, 0.00522619066759944f, 0.004349356517195702f, 0.00330205773934722f, 0.002286989474669099f, 0.001444985857233405f, 0.0008328807889483869f, 0.0004379475140012801f, 0.0002100782585330308f, 9.193058212986216e-005f,\n"
"5.808438800158911e-005f, 0.0001327334903180599f, 0.0002767078403849155f, 0.0005262380582280457f, 0.0009129836107604206f, 0.001444985857233405f, 0.002086335094645619f, 0.002748048631474376f, 0.00330205773934722f, 0.003619635012000799f, 0.003619635012000799f, 0.00330205773934722f, 0.002748048631474376f, 0.002086335094645619f, 0.001444985857233405f, 0.0009129836107604206f, 0.0005262380582280457f, 0.0002767078403849155f, 0.0001327334903180599f, 5.808438800158911e-005f,\n"
"3.34794785885606e-005f, 7.650675252079964e-005f, 0.0001594926579855382f, 0.0003033203247468919f, 0.0005262380582280457f, 0.0008328807889483869f, 0.001202550483867526f, 0.001583957928232849f, 0.001903285388834775f, 0.002086334861814976f, 0.002086334861814976f, 0.001903285388834775f, 0.001583957928232849f, 0.001202550483867526f, 0.0008328807889483869f, 0.0005262380582280457f, 0.0003033203247468919f, 0.0001594926579855382f, 7.650675252079964e-005f, 3.34794785885606e-005f,\n"
"1.760426494001877e-005f, 4.022897701361217e-005f, 8.386484114453197e-005f, 0.0001594926579855382f, 0.0002767078403849155f, 0.0004379475140012801f, 0.0006323281559161842f, 0.0008328808471560478f, 0.001000790391117334f, 0.001097041997127235f, 0.001097041997127235f, 0.001000790391117334f, 0.0008328808471560478f, 0.0006323281559161842f, 0.0004379475140012801f, 0.0002767078403849155f, 0.0001594926579855382f, 8.386484114453197e-005f, 4.022897701361217e-005f, 1.760426494001877e-005f,\n"
"8.444558261544444e-006f, 1.929736572492402e-005f, 4.022897701361217e-005f, 7.650675252079964e-005f, 0.0001327334903180599f, 0.0002100782585330308f, 0.0003033203829545528f, 0.0003995231236331165f, 0.0004800673632416874f, 0.0005262381164357066f, 0.0005262381164357066f, 0.0004800673632416874f, 0.0003995231236331165f, 0.0003033203829545528f, 0.0002100782585330308f, 0.0001327334903180599f, 7.650675252079964e-005f, 4.022897701361217e-005f, 1.929736572492402e-005f, 8.444558261544444e-006f,\n"
"3.695352233989979e-006f, 8.444558261544444e-006f, 1.760426494001877e-005f, 3.34794785885606e-005f, 5.808438800158911e-005f, 9.193058212986216e-005f, 0.0001327334757661447f, 0.0001748319627949968f, 0.0002100782439811155f, 0.0002302826324012131f, 0.0002302826324012131f, 0.0002100782439811155f, 0.0001748319627949968f, 0.0001327334757661447f, 9.193058212986216e-005f, 5.808438800158911e-005f, 3.34794785885606e-005f, 1.760426494001877e-005f, 8.444558261544444e-006f, 3.695352233989979e-006f\n"
"};\n"
"#define readerGet(centerX, centerY, win_offset, cos_dir, sin_dir, i, j) \\\n"
"read_imgTex((float2)(centerX + (win_offset + j) * cos_dir + (win_offset + i) * sin_dir, \\\n"
"centerY - (win_offset + j) * sin_dir + (win_offset + i) * cos_dir))\n"
"inline float linearFilter(\n"
"__PARAM_imgTex__, int img_rows, int img_cols,\n"
"float centerX, float centerY, float win_offset,\n"
"float cos_dir, float sin_dir, float y, float x )\n"
"{\n"
"float out = 0.0f;\n"
"const int x1 = round(x);\n"
"const int y1 = round(y);\n"
"const int x2 = x1 + 1;\n"
"const int y2 = y1 + 1;\n"
"uchar src_reg = readerGet(centerX, centerY, win_offset, cos_dir, sin_dir, y1, x1);\n"
"out = out + src_reg * ((x2 - x) * (y2 - y));\n"
"src_reg = readerGet(centerX, centerY, win_offset, cos_dir, sin_dir, y1, x2);\n"
"out = out + src_reg * ((x - x1) * (y2 - y));\n"
"src_reg = readerGet(centerX, centerY, win_offset, cos_dir, sin_dir, y2, x1);\n"
"out = out + src_reg * ((x2 - x) * (y - y1));\n"
"src_reg = readerGet(centerX, centerY, win_offset, cos_dir, sin_dir, y2, x2);\n"
"out = out + src_reg * ((x - x1) * (y - y1));\n"
"return out;\n"
"}\n"
"inline float areaFilter( __PARAM_imgTex__, int img_rows, int img_cols,\n"
"float centerX, float centerY, float win_offset,\n"
"float cos_dir, float sin_dir, float x, float y, float s)\n"
"{\n"
"float fsx1 = x * s;\n"
"float fsx2 = fsx1 + s;\n"
"int sx1 = convert_int_rtp(fsx1);\n"
"int sx2 = convert_int_rtn(fsx2);\n"
"float fsy1 = y * s;\n"
"float fsy2 = fsy1 + s;\n"
"int sy1 = convert_int_rtp(fsy1);\n"
"int sy2 = convert_int_rtn(fsy2);\n"
"float scale = 1.f / (s * s);\n"
"float out = 0.f;\n"
"for (int dy = sy1; dy < sy2; ++dy)\n"
"{\n"
"for (int dx = sx1; dx < sx2; ++dx)\n"
"out = out + readerGet(centerX, centerY, win_offset, cos_dir, sin_dir, dy, dx) * scale;\n"
"if (sx1 > fsx1)\n"
"out = out + readerGet(centerX, centerY, win_offset, cos_dir, sin_dir, dy, (sx1 -1)) * ((sx1 - fsx1) * scale);\n"
"if (sx2 < fsx2)\n"
"out = out + readerGet(centerX, centerY, win_offset, cos_dir, sin_dir, dy, sx2) * ((fsx2 -sx2) * scale);\n"
"}\n"
"if (sy1 > fsy1)\n"
"for (int dx = sx1; dx < sx2; ++dx)\n"
"out = out + readerGet(centerX, centerY, win_offset, cos_dir, sin_dir, (sy1 - 1) , dx) * ((sy1 -fsy1) * scale);\n"
"if (sy2 < fsy2)\n"
"for (int dx = sx1; dx < sx2; ++dx)\n"
"out = out + readerGet(centerX, centerY, win_offset, cos_dir, sin_dir, sy2, dx) * ((fsy2 -sy2) * scale);\n"
"if ((sy1 > fsy1) &&  (sx1 > fsx1))\n"
"out = out + readerGet(centerX, centerY, win_offset, cos_dir, sin_dir, (sy1 - 1) , (sx1 - 1)) * ((sy1 -fsy1) * (sx1 -fsx1) * scale);\n"
"if ((sy1 > fsy1) &&  (sx2 < fsx2))\n"
"out = out + readerGet(centerX, centerY, win_offset, cos_dir, sin_dir, (sy1 - 1) , sx2) * ((sy1 -fsy1) * (fsx2 -sx2) * scale);\n"
"if ((sy2 < fsy2) &&  (sx2 < fsx2))\n"
"out = out + readerGet(centerX, centerY, win_offset, cos_dir, sin_dir, sy2, sx2) * ((fsy2 -sy2) * (fsx2 -sx2) * scale);\n"
"if ((sy2 < fsy2) &&  (sx1 > fsx1))\n"
"out = out + readerGet(centerX, centerY, win_offset, cos_dir, sin_dir, sy2, (sx1 - 1)) * ((fsy2 -sy2) * (sx1 -fsx1) * scale);\n"
"return out;\n"
"}\n"
"void calc_dx_dy(\n"
"__PARAM_imgTex__,\n"
"int img_rows, int img_cols,\n"
"volatile __local  float *s_dx_bin,\n"
"volatile __local  float *s_dy_bin,\n"
"volatile __local  float *s_PATCH,\n"
"__global const float* featureX,\n"
"__global const float* featureY,\n"
"__global const float* featureSize,\n"
"__global const float* featureDir )\n"
"{\n"
"const float centerX = featureX[get_group_id(0)];\n"
"const float centerY = featureY[get_group_id(0)];\n"
"const float size = featureSize[get_group_id(0)];\n"
"float descriptor_dir = 360.0f - featureDir[get_group_id(0)];\n"
"if(fabs(descriptor_dir - 360.0f) < FLT_EPSILON)\n"
"{\n"
"descriptor_dir = 0.0f;\n"
"}\n"
"descriptor_dir *= (float)(CV_PI_F / 180.0f);\n"
"const float s = size * 1.2f / 9.0f;\n"
"const int win_size = (int)((PATCH_SZ + 1) * s);\n"
"float sin_dir;\n"
"float cos_dir;\n"
"sin_dir = sincos(descriptor_dir, &cos_dir);\n"
"const float win_offset = -(float)(win_size - 1) / 2;\n"
"const int xBlock = (get_group_id(1) & 3);\n"
"const int yBlock = (get_group_id(1) >> 2);\n"
"const int xIndex = xBlock * 5 + get_local_id(0);\n"
"const int yIndex = yBlock * 5 + get_local_id(1);\n"
"const float icoo = ((float)yIndex / (PATCH_SZ + 1)) * win_size;\n"
"const float jcoo = ((float)xIndex / (PATCH_SZ + 1)) * win_size;\n"
"if (s > 1)\n"
"{\n"
"s_PATCH[get_local_id(1) * 6 + get_local_id(0)] =\n"
"areaFilter(__PASS_imgTex__, img_rows, img_cols, centerX, centerY,\n"
"win_offset, cos_dir, sin_dir, xIndex, yIndex, s);\n"
"}\n"
"else\n"
"{\n"
"s_PATCH[get_local_id(1) * 6 + get_local_id(0)] =\n"
"linearFilter(__PASS_imgTex__, img_rows, img_cols, centerX, centerY,\n"
"win_offset, cos_dir, sin_dir, icoo, jcoo);\n"
"}\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"if (get_local_id(0) < 5 && get_local_id(1) < 5)\n"
"{\n"
"const int tid = get_local_id(1) * 5 + get_local_id(0);\n"
"const float dw = c_DW[yIndex * PATCH_SZ + xIndex];\n"
"const float vx = (\n"
"s_PATCH[      get_local_id(1) * 6 + get_local_id(0) + 1] -\n"
"s_PATCH[      get_local_id(1) * 6 + get_local_id(0)    ] +\n"
"s_PATCH[(get_local_id(1) + 1) * 6 + get_local_id(0) + 1] -\n"
"s_PATCH[(get_local_id(1) + 1) * 6 + get_local_id(0)    ])\n"
"* dw;\n"
"const float vy = (\n"
"s_PATCH[(get_local_id(1) + 1) * 6 + get_local_id(0)    ] -\n"
"s_PATCH[      get_local_id(1) * 6 + get_local_id(0)    ] +\n"
"s_PATCH[(get_local_id(1) + 1) * 6 + get_local_id(0) + 1] -\n"
"s_PATCH[      get_local_id(1) * 6 + get_local_id(0) + 1])\n"
"* dw;\n"
"s_dx_bin[tid] = vx;\n"
"s_dy_bin[tid] = vy;\n"
"}\n"
"}\n"
"void reduce_sum25(\n"
"volatile __local  float* sdata1,\n"
"volatile __local  float* sdata2,\n"
"volatile __local  float* sdata3,\n"
"volatile __local  float* sdata4,\n"
"int tid\n"
")\n"
"{\n"
"#ifndef WAVE_SIZE\n"
"#define WAVE_SIZE 1\n"
"#endif\n"
"if (tid < 9)\n"
"{\n"
"sdata1[tid] += sdata1[tid + 16];\n"
"sdata2[tid] += sdata2[tid + 16];\n"
"sdata3[tid] += sdata3[tid + 16];\n"
"sdata4[tid] += sdata4[tid + 16];\n"
"#if WAVE_SIZE < 16\n"
"}\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"if (tid < 8)\n"
"{\n"
"#endif\n"
"sdata1[tid] += sdata1[tid + 8];\n"
"sdata2[tid] += sdata2[tid + 8];\n"
"sdata3[tid] += sdata3[tid + 8];\n"
"sdata4[tid] += sdata4[tid + 8];\n"
"#if WAVE_SIZE < 8\n"
"}\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"if (tid < 4)\n"
"{\n"
"#endif\n"
"sdata1[tid] += sdata1[tid + 4];\n"
"sdata2[tid] += sdata2[tid + 4];\n"
"sdata3[tid] += sdata3[tid + 4];\n"
"sdata4[tid] += sdata4[tid + 4];\n"
"#if WAVE_SIZE < 4\n"
"}\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"if (tid < 2)\n"
"{\n"
"#endif\n"
"sdata1[tid] += sdata1[tid + 2];\n"
"sdata2[tid] += sdata2[tid + 2];\n"
"sdata3[tid] += sdata3[tid + 2];\n"
"sdata4[tid] += sdata4[tid + 2];\n"
"#if WAVE_SIZE < 2\n"
"}\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"if (tid < 1)\n"
"{\n"
"#endif\n"
"sdata1[tid] += sdata1[tid + 1];\n"
"sdata2[tid] += sdata2[tid + 1];\n"
"sdata3[tid] += sdata3[tid + 1];\n"
"sdata4[tid] += sdata4[tid + 1];\n"
"}\n"
"#undef WAVE_SIZE\n"
"}\n"
"__kernel\n"
"void SURF_computeDescriptors64(\n"
"__PARAM_imgTex__,\n"
"int img_rows, int img_cols,\n"
"__global const float* keypoints,\n"
"int keypoints_step, int keypoints_offset,\n"
"__global float * descriptors,\n"
"int descriptors_step, int descriptors_offset)\n"
"{\n"
"descriptors_step /= sizeof(float);\n"
"keypoints_step   /= sizeof(float);\n"
"__global const float * featureX    = keypoints + X_ROW * keypoints_step;\n"
"__global const float * featureY    = keypoints + Y_ROW * keypoints_step;\n"
"__global const float * featureSize = keypoints + SIZE_ROW * keypoints_step;\n"
"__global const float * featureDir  = keypoints + ANGLE_ROW * keypoints_step;\n"
"volatile __local  float sdx[25];\n"
"volatile __local  float sdy[25];\n"
"volatile __local  float sdxabs[25];\n"
"volatile __local  float sdyabs[25];\n"
"volatile __local  float s_PATCH[6*6];\n"
"calc_dx_dy(__PASS_imgTex__, img_rows, img_cols, sdx, sdy, s_PATCH, featureX, featureY, featureSize, featureDir);\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"const int tid = get_local_id(1) * get_local_size(0) + get_local_id(0);\n"
"if (tid < 25)\n"
"{\n"
"sdxabs[tid] = fabs(sdx[tid]);\n"
"sdyabs[tid] = fabs(sdy[tid]);\n"
"}\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"reduce_sum25(sdx, sdy, sdxabs, sdyabs, tid);\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"if (tid == 0)\n"
"{\n"
"__global float* descriptors_block = descriptors + descriptors_step * get_group_id(0) + (get_group_id(1) << 2);\n"
"descriptors_block[0] = sdx[0];\n"
"descriptors_block[1] = sdy[0];\n"
"descriptors_block[2] = sdxabs[0];\n"
"descriptors_block[3] = sdyabs[0];\n"
"}\n"
"}\n"
"__kernel\n"
"void SURF_computeDescriptors128(\n"
"__PARAM_imgTex__,\n"
"int img_rows, int img_cols,\n"
"__global const float* keypoints,\n"
"int keypoints_step, int keypoints_offset,\n"
"__global float* descriptors,\n"
"int descriptors_step, int descriptors_offset)\n"
"{\n"
"descriptors_step /= sizeof(*descriptors);\n"
"keypoints_step   /= sizeof(*keypoints);\n"
"__global const float * featureX   = keypoints + X_ROW * keypoints_step;\n"
"__global const float * featureY   = keypoints + Y_ROW * keypoints_step;\n"
"__global const float* featureSize = keypoints + SIZE_ROW * keypoints_step;\n"
"__global const float* featureDir  = keypoints + ANGLE_ROW * keypoints_step;\n"
"volatile __local  float sdx[25];\n"
"volatile __local  float sdy[25];\n"
"volatile __local  float sd1[25];\n"
"volatile __local  float sd2[25];\n"
"volatile __local  float sdabs1[25];\n"
"volatile __local  float sdabs2[25];\n"
"volatile __local  float s_PATCH[6*6];\n"
"calc_dx_dy(__PASS_imgTex__, img_rows, img_cols, sdx, sdy, s_PATCH, featureX, featureY, featureSize, featureDir);\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"const int tid = get_local_id(1) * get_local_size(0) + get_local_id(0);\n"
"if (tid < 25)\n"
"{\n"
"if (sdy[tid] >= 0)\n"
"{\n"
"sd1[tid] = sdx[tid];\n"
"sdabs1[tid] = fabs(sdx[tid]);\n"
"sd2[tid] = 0;\n"
"sdabs2[tid] = 0;\n"
"}\n"
"else\n"
"{\n"
"sd1[tid] = 0;\n"
"sdabs1[tid] = 0;\n"
"sd2[tid] = sdx[tid];\n"
"sdabs2[tid] = fabs(sdx[tid]);\n"
"}\n"
"}\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"reduce_sum25(sd1, sd2, sdabs1, sdabs2, tid);\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"__global float* descriptors_block = descriptors + descriptors_step * get_group_id(0) + (get_group_id(1) << 3);\n"
"if (tid < 25)\n"
"{\n"
"if (tid == 0)\n"
"{\n"
"descriptors_block[0] = sd1[0];\n"
"descriptors_block[1] = sdabs1[0];\n"
"descriptors_block[2] = sd2[0];\n"
"descriptors_block[3] = sdabs2[0];\n"
"}\n"
"if (sdx[tid] >= 0)\n"
"{\n"
"sd1[tid] = sdy[tid];\n"
"sdabs1[tid] = fabs(sdy[tid]);\n"
"sd2[tid] = 0;\n"
"sdabs2[tid] = 0;\n"
"}\n"
"else\n"
"{\n"
"sd1[tid] = 0;\n"
"sdabs1[tid] = 0;\n"
"sd2[tid] = sdy[tid];\n"
"sdabs2[tid] = fabs(sdy[tid]);\n"
"}\n"
"}\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"reduce_sum25(sd1, sd2, sdabs1, sdabs2, tid);\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"if (tid < 25)\n"
"{\n"
"if (tid == 0)\n"
"{\n"
"descriptors_block[4] = sd1[0];\n"
"descriptors_block[5] = sdabs1[0];\n"
"descriptors_block[6] = sd2[0];\n"
"descriptors_block[7] = sdabs2[0];\n"
"}\n"
"}\n"
"}\n"
"void reduce_sum128(volatile __local  float* smem, int tid)\n"
"{\n"
"#ifndef WAVE_SIZE\n"
"#define WAVE_SIZE 1\n"
"#endif\n"
"if (tid < 64)\n"
"{\n"
"smem[tid] += smem[tid + 64];\n"
"#if WAVE_SIZE < 64\n"
"}\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"if (tid < 32)\n"
"{\n"
"#endif\n"
"smem[tid] += smem[tid + 32];\n"
"#if WAVE_SIZE < 32\n"
"}\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"if (tid < 16)\n"
"{\n"
"#endif\n"
"smem[tid] += smem[tid + 16];\n"
"#if WAVE_SIZE < 16\n"
"}\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"if (tid < 8)\n"
"{\n"
"#endif\n"
"smem[tid] += smem[tid + 8];\n"
"#if WAVE_SIZE < 8\n"
"}\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"if (tid < 4)\n"
"{\n"
"#endif\n"
"smem[tid] += smem[tid + 4];\n"
"#if WAVE_SIZE < 4\n"
"}\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"if (tid < 2)\n"
"{\n"
"#endif\n"
"smem[tid] += smem[tid + 2];\n"
"#if WAVE_SIZE < 2\n"
"}\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"if (tid < 1)\n"
"{\n"
"#endif\n"
"smem[tid] += smem[tid + 1];\n"
"}\n"
"}\n"
"void reduce_sum64(volatile __local  float* smem, int tid)\n"
"{\n"
"#ifndef WAVE_SIZE\n"
"#define WAVE_SIZE 1\n"
"#endif\n"
"if (tid < 32)\n"
"{\n"
"smem[tid] += smem[tid + 32];\n"
"#if WAVE_SIZE < 32\n"
"}\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"if (tid < 16)\n"
"{\n"
"#endif\n"
"smem[tid] += smem[tid + 16];\n"
"#if WAVE_SIZE < 16\n"
"}\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"if (tid < 8)\n"
"{\n"
"#endif\n"
"smem[tid] += smem[tid + 8];\n"
"#if WAVE_SIZE < 8\n"
"}\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"if (tid < 4)\n"
"{\n"
"#endif\n"
"smem[tid] += smem[tid + 4];\n"
"#if WAVE_SIZE < 4\n"
"}\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"if (tid < 2)\n"
"{\n"
"#endif\n"
"smem[tid] += smem[tid + 2];\n"
"#if WAVE_SIZE < 2\n"
"}\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"if (tid < 1)\n"
"{\n"
"#endif\n"
"smem[tid] += smem[tid + 1];\n"
"}\n"
"}\n"
"__kernel\n"
"void SURF_normalizeDescriptors128(__global float * descriptors, int descriptors_step, int descriptors_offset)\n"
"{\n"
"descriptors_step /= sizeof(*descriptors);\n"
"__global float* descriptor_base = descriptors + descriptors_step * get_group_id(0);\n"
"volatile __local  float sqDesc[128];\n"
"const float lookup = descriptor_base[get_local_id(0)];\n"
"sqDesc[get_local_id(0)] = lookup * lookup;\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"reduce_sum128(sqDesc, get_local_id(0));\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"volatile __local  float len;\n"
"if (get_local_id(0) == 0)\n"
"{\n"
"len = sqrt(sqDesc[0]);\n"
"}\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"descriptor_base[get_local_id(0)] = lookup / len;\n"
"}\n"
"__kernel\n"
"void SURF_normalizeDescriptors64(__global float * descriptors, int descriptors_step, int descriptors_offset)\n"
"{\n"
"descriptors_step /= sizeof(*descriptors);\n"
"__global float* descriptor_base = descriptors + descriptors_step * get_group_id(0);\n"
"volatile __local  float sqDesc[64];\n"
"const float lookup = descriptor_base[get_local_id(0)];\n"
"sqDesc[get_local_id(0)] = lookup * lookup;\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"reduce_sum64(sqDesc, get_local_id(0));\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"volatile __local  float len;\n"
"if (get_local_id(0) == 0)\n"
"{\n"
"len = sqrt(sqDesc[0]);\n"
"}\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"descriptor_base[get_local_id(0)] = lookup / len;\n"
"}\n"
, "fa320902fb1b95488f0dc35a0fa12041", NULL};

}}}
#endif