1 // This file is auto-generated. Do not edit! 2 3 #include "precomp.hpp" 4 #include "opencl_kernels_stitching.hpp" 5 6 namespace cv 7 { 8 namespace ocl 9 { 10 namespace stitching 11 { 12 13 const struct ProgramEntry multibandblend={"multibandblend", 14 "#ifndef NL\n" 15 "#define NL\n" 16 "#endif\n" 17 "#define REF(x) x\n" 18 "#define __CAT(x, y) x##y\n" 19 "#define CAT(x, y) __CAT(x, y)\n" 20 "#define DECLARE_MAT_ARG(name) \\\n" 21 "__global uchar* restrict name ## Ptr, \\\n" 22 "int name ## StepBytes, \\\n" 23 "int name ## Offset, \\\n" 24 "int name ## Height, \\\n" 25 "int name ## Width NL\n" 26 "#define MAT_BYTE_OFFSET(name, x, y) mad24((y), name ## StepBytes, ((x)) * (int)(name ## _TSIZE) + name ## Offset)\n" 27 "#define MAT_RELATIVE_BYTE_OFFSET(name, x, y) mad24(y, name ## StepBytes, (x) * (int)(name ## _TSIZE))\n" 28 "#define __LOAD_MAT_AT(name, byteOffset) *((const __global name ## _T*)(name ## Ptr + (byteOffset)))\n" 29 "#define __vload_CN__(name_cn) vload ## name_cn\n" 30 "#define __vload_CN_(name_cn) __vload_CN__(name_cn)\n" 31 "#define __vload_CN(name) __vload_CN_(name ## _CN)\n" 32 "#define __LOAD_MAT_AT_vload(name, byteOffset) __vload_CN(name)(0, ((const __global name ## _T1*)(name ## Ptr + (byteOffset))))\n" 33 "#define __LOAD_MAT_AT_1 __LOAD_MAT_AT\n" 34 "#define __LOAD_MAT_AT_2 __LOAD_MAT_AT\n" 35 "#define __LOAD_MAT_AT_3 __LOAD_MAT_AT_vload\n" 36 "#define __LOAD_MAT_AT_4 __LOAD_MAT_AT\n" 37 "#define __LOAD_MAT_AT_CN__(name_cn) __LOAD_MAT_AT_ ## name_cn\n" 38 "#define __LOAD_MAT_AT_CN_(name_cn) __LOAD_MAT_AT_CN__(name_cn)\n" 39 "#define __LOAD_MAT_AT_CN(name) __LOAD_MAT_AT_CN_(name ## _CN)\n" 40 "#define LOAD_MAT_AT(name, byteOffset) __LOAD_MAT_AT_CN(name)(name, byteOffset)\n" 41 "#define __STORE_MAT_AT(name, byteOffset, v) *((__global name ## _T*)(name ## Ptr + (byteOffset))) = v\n" 42 "#define __vstore_CN__(name_cn) vstore ## name_cn\n" 43 "#define __vstore_CN_(name_cn) __vstore_CN__(name_cn)\n" 44 "#define __vstore_CN(name) __vstore_CN_(name ## _CN)\n" 45 "#define __STORE_MAT_AT_vstore(name, byteOffset, v) __vstore_CN(name)(v, 0, ((__global name ## _T1*)(name ## Ptr + (byteOffset))))\n" 46 "#define __STORE_MAT_AT_1 __STORE_MAT_AT\n" 47 "#define __STORE_MAT_AT_2 __STORE_MAT_AT\n" 48 "#define __STORE_MAT_AT_3 __STORE_MAT_AT_vstore\n" 49 "#define __STORE_MAT_AT_4 __STORE_MAT_AT\n" 50 "#define __STORE_MAT_AT_CN__(name_cn) __STORE_MAT_AT_ ## name_cn\n" 51 "#define __STORE_MAT_AT_CN_(name_cn) __STORE_MAT_AT_CN__(name_cn)\n" 52 "#define __STORE_MAT_AT_CN(name) __STORE_MAT_AT_CN_(name ## _CN)\n" 53 "#define STORE_MAT_AT(name, byteOffset, v) __STORE_MAT_AT_CN(name)(name, byteOffset, v)\n" 54 "#define T1_uchar uchar\n" 55 "#define T1_uchar2 uchar\n" 56 "#define T1_uchar3 uchar\n" 57 "#define T1_uchar4 uchar\n" 58 "#define T1_char char\n" 59 "#define T1_char2 char\n" 60 "#define T1_char3 char\n" 61 "#define T1_char4 char\n" 62 "#define T1_ushort ushort\n" 63 "#define T1_ushort2 ushort\n" 64 "#define T1_ushort3 ushort\n" 65 "#define T1_ushort4 ushort\n" 66 "#define T1_short short\n" 67 "#define T1_short2 short\n" 68 "#define T1_short3 short\n" 69 "#define T1_short4 short\n" 70 "#define T1_int int\n" 71 "#define T1_int2 int\n" 72 "#define T1_int3 int\n" 73 "#define T1_int4 int\n" 74 "#define T1_float float\n" 75 "#define T1_float2 float\n" 76 "#define T1_float3 float\n" 77 "#define T1_float4 float\n" 78 "#define T1_double double\n" 79 "#define T1_double2 double\n" 80 "#define T1_double3 double\n" 81 "#define T1_double4 double\n" 82 "#define T1(type) REF(CAT(T1_, REF(type)))\n" 83 "#define uchar1 uchar\n" 84 "#define char1 char\n" 85 "#define short1 short\n" 86 "#define ushort1 ushort\n" 87 "#define int1 int\n" 88 "#define float1 float\n" 89 "#define double1 double\n" 90 "#define TYPE(type, cn) REF(CAT(REF(type), REF(cn)))\n" 91 "#define __CONVERT_MODE_uchar_uchar __NO_CONVERT\n" 92 "#define __CONVERT_MODE_uchar_char __CONVERT_sat\n" 93 "#define __CONVERT_MODE_uchar_ushort __CONVERT\n" 94 "#define __CONVERT_MODE_uchar_short __CONVERT\n" 95 "#define __CONVERT_MODE_uchar_int __CONVERT\n" 96 "#define __CONVERT_MODE_uchar_float __CONVERT\n" 97 "#define __CONVERT_MODE_uchar_double __CONVERT\n" 98 "#define __CONVERT_MODE_char_uchar __CONVERT_sat\n" 99 "#define __CONVERT_MODE_char_char __NO_CONVERT\n" 100 "#define __CONVERT_MODE_char_ushort __CONVERT_sat\n" 101 "#define __CONVERT_MODE_char_short __CONVERT\n" 102 "#define __CONVERT_MODE_char_int __CONVERT\n" 103 "#define __CONVERT_MODE_char_float __CONVERT\n" 104 "#define __CONVERT_MODE_char_double __CONVERT\n" 105 "#define __CONVERT_MODE_ushort_uchar __CONVERT_sat\n" 106 "#define __CONVERT_MODE_ushort_char __CONVERT_sat\n" 107 "#define __CONVERT_MODE_ushort_ushort __NO_CONVERT\n" 108 "#define __CONVERT_MODE_ushort_short __CONVERT_sat\n" 109 "#define __CONVERT_MODE_ushort_int __CONVERT\n" 110 "#define __CONVERT_MODE_ushort_float __CONVERT\n" 111 "#define __CONVERT_MODE_ushort_double __CONVERT\n" 112 "#define __CONVERT_MODE_short_uchar __CONVERT_sat\n" 113 "#define __CONVERT_MODE_short_char __CONVERT_sat\n" 114 "#define __CONVERT_MODE_short_ushort __CONVERT_sat\n" 115 "#define __CONVERT_MODE_short_short __NO_CONVERT\n" 116 "#define __CONVERT_MODE_short_int __CONVERT\n" 117 "#define __CONVERT_MODE_short_float __CONVERT\n" 118 "#define __CONVERT_MODE_short_double __CONVERT\n" 119 "#define __CONVERT_MODE_int_uchar __CONVERT_sat\n" 120 "#define __CONVERT_MODE_int_char __CONVERT_sat\n" 121 "#define __CONVERT_MODE_int_ushort __CONVERT_sat\n" 122 "#define __CONVERT_MODE_int_short __CONVERT_sat\n" 123 "#define __CONVERT_MODE_int_int __NO_CONVERT\n" 124 "#define __CONVERT_MODE_int_float __CONVERT\n" 125 "#define __CONVERT_MODE_int_double __CONVERT\n" 126 "#define __CONVERT_MODE_float_uchar __CONVERT_sat_rte\n" 127 "#define __CONVERT_MODE_float_char __CONVERT_sat_rte\n" 128 "#define __CONVERT_MODE_float_ushort __CONVERT_sat_rte\n" 129 "#define __CONVERT_MODE_float_short __CONVERT_sat_rte\n" 130 "#define __CONVERT_MODE_float_int __CONVERT_rte\n" 131 "#define __CONVERT_MODE_float_float __NO_CONVERT\n" 132 "#define __CONVERT_MODE_float_double __CONVERT\n" 133 "#define __CONVERT_MODE_double_uchar __CONVERT_sat_rte\n" 134 "#define __CONVERT_MODE_double_char __CONVERT_sat_rte\n" 135 "#define __CONVERT_MODE_double_ushort __CONVERT_sat_rte\n" 136 "#define __CONVERT_MODE_double_short __CONVERT_sat_rte\n" 137 "#define __CONVERT_MODE_double_int __CONVERT_rte\n" 138 "#define __CONVERT_MODE_double_float __CONVERT\n" 139 "#define __CONVERT_MODE_double_double __NO_CONVERT\n" 140 "#define __CONVERT_MODE(srcType, dstType) CAT(__CONVERT_MODE_, CAT(REF(T1(srcType)), CAT(_, REF(T1(dstType)))))\n" 141 "#define __ROUND_MODE__NO_CONVERT\n" 142 "#define __ROUND_MODE__CONVERT\n" 143 "#define __ROUND_MODE__CONVERT_rte _rte\n" 144 "#define __ROUND_MODE__CONVERT_sat _sat\n" 145 "#define __ROUND_MODE__CONVERT_sat_rte _sat_rte\n" 146 "#define ROUND_MODE(srcType, dstType) CAT(__ROUND_MODE_, __CONVERT_MODE(srcType, dstType))\n" 147 "#define __CONVERT_ROUND(dstType, roundMode) CAT(CAT(convert_, REF(dstType)), roundMode)\n" 148 "#define __NO_CONVERT(dstType)\n" 149 "#define __CONVERT(dstType) __CONVERT_ROUND(dstType,)\n" 150 "#define __CONVERT_rte(dstType) __CONVERT_ROUND(dstType,_rte)\n" 151 "#define __CONVERT_sat(dstType) __CONVERT_ROUND(dstType,_sat)\n" 152 "#define __CONVERT_sat_rte(dstType) __CONVERT_ROUND(dstType,_sat_rte)\n" 153 "#define CONVERT(srcType, dstType) REF(__CONVERT_MODE(srcType,dstType))(dstType)\n" 154 "#define CONVERT_TO(dstType) __CONVERT_ROUND(dstType,)\n" 155 "#define CV_8U 0\n" 156 "#define CV_8S 1\n" 157 "#define CV_16U 2\n" 158 "#define CV_16S 3\n" 159 "#define CV_32S 4\n" 160 "#define CV_32F 5\n" 161 "#define CV_64F 6\n" 162 "#if defined(DEFINE_feed)\n" 163 "#define workType TYPE(weight_T1, src_CN)\n" 164 "#if src_DEPTH == 3 && src_CN == 3\n" 165 "#define convertSrcToWorkType convert_float3\n" 166 "#else\n" 167 "#define convertSrcToWorkType CONVERT_TO(workType)\n" 168 "#endif\n" 169 "#if dst_DEPTH == 3 && dst_CN == 3\n" 170 "#define convertToDstType convert_short3\n" 171 "#else\n" 172 "#define convertToDstType CONVERT_TO(dst_T)\n" 173 "#endif\n" 174 "__kernel void feed(\n" 175 "DECLARE_MAT_ARG(src), DECLARE_MAT_ARG(weight),\n" 176 "DECLARE_MAT_ARG(dst), DECLARE_MAT_ARG(dstWeight)\n" 177 ")\n" 178 "{\n" 179 "const int x = get_global_id(0);\n" 180 "const int y = get_global_id(1);\n" 181 "if (x < srcWidth && y < srcHeight)\n" 182 "{\n" 183 "int src_byteOffset = MAT_BYTE_OFFSET(src, x, y);\n" 184 "int weight_byteOffset = MAT_BYTE_OFFSET(weight, x, y);\n" 185 "int dst_byteOffset = MAT_BYTE_OFFSET(dst, x, y);\n" 186 "int dstWeight_byteOffset = MAT_BYTE_OFFSET(dstWeight, x, y);\n" 187 "weight_T w = LOAD_MAT_AT(weight, weight_byteOffset);\n" 188 "workType src_value = convertSrcToWorkType(LOAD_MAT_AT(src, src_byteOffset));\n" 189 "STORE_MAT_AT(dst, dst_byteOffset, LOAD_MAT_AT(dst, dst_byteOffset) + convertToDstType(src_value * w));\n" 190 "STORE_MAT_AT(dstWeight, dstWeight_byteOffset, LOAD_MAT_AT(dstWeight, dstWeight_byteOffset) + w);\n" 191 "}\n" 192 "}\n" 193 "#endif\n" 194 "#if defined(DEFINE_normalizeUsingWeightMap)\n" 195 "#if mat_DEPTH == 3 && mat_CN == 3\n" 196 "#define workType float3\n" 197 "#define convertSrcToWorkType convert_float3\n" 198 "#define convertToDstType convert_short3\n" 199 "#else\n" 200 "#define workType TYPE(weight_T1, mat_CN)\n" 201 "#define convertSrcToWorkType CONVERT_TO(workType)\n" 202 "#define convertToDstType CONVERT_TO(mat_T)\n" 203 "#endif\n" 204 "#if weight_DEPTH >= CV_32F\n" 205 "#define WEIGHT_EPS 1e-5f\n" 206 "#else\n" 207 "#define WEIGHT_EPS 0\n" 208 "#endif\n" 209 "__kernel void normalizeUsingWeightMap(\n" 210 "DECLARE_MAT_ARG(mat), DECLARE_MAT_ARG(weight)\n" 211 ")\n" 212 "{\n" 213 "const int x = get_global_id(0);\n" 214 "const int y = get_global_id(1);\n" 215 "if (x < matWidth && y < matHeight)\n" 216 "{\n" 217 "int mat_byteOffset = MAT_BYTE_OFFSET(mat, x, y);\n" 218 "int weight_byteOffset = MAT_BYTE_OFFSET(weight, x, y);\n" 219 "weight_T w = LOAD_MAT_AT(weight, weight_byteOffset);\n" 220 "workType value = convertSrcToWorkType(LOAD_MAT_AT(mat, mat_byteOffset));\n" 221 "value = value / (w + WEIGHT_EPS);\n" 222 "STORE_MAT_AT(mat, mat_byteOffset, convertToDstType(value));\n" 223 "}\n" 224 "}\n" 225 "#endif\n" 226 , "3320d5f13a357c8ee3c223e66d598244"}; 227 ProgramSource multibandblend_oclsrc(multibandblend.programStr); 228 const struct ProgramEntry warpers={"warpers", 229 "__kernel void buildWarpPlaneMaps(__global uchar * xmapptr, int xmap_step, int xmap_offset,\n" 230 "__global uchar * ymapptr, int ymap_step, int ymap_offset, int rows, int cols,\n" 231 "__constant float * ck_rinv, __constant float * ct,\n" 232 "int tl_u, int tl_v, float scale, int rowsPerWI)\n" 233 "{\n" 234 "int du = get_global_id(0);\n" 235 "int dv0 = get_global_id(1) * rowsPerWI;\n" 236 "if (du < cols)\n" 237 "{\n" 238 "int xmap_index = mad24(dv0, xmap_step, mad24(du, (int)sizeof(float), xmap_offset));\n" 239 "int ymap_index = mad24(dv0, ymap_step, mad24(du, (int)sizeof(float), ymap_offset));\n" 240 "float u = tl_u + du;\n" 241 "float x_ = fma(u, scale, -ct[0]);\n" 242 "float ct1 = 1 - ct[2];\n" 243 "for (int dv = dv0, dv1 = min(rows, dv0 + rowsPerWI); dv < dv1; ++dv, xmap_index += xmap_step,\n" 244 "ymap_index += ymap_step)\n" 245 "{\n" 246 "__global float * xmap = (__global float *)(xmapptr + xmap_index);\n" 247 "__global float * ymap = (__global float *)(ymapptr + ymap_index);\n" 248 "float v = tl_v + dv;\n" 249 "float y_ = fma(v, scale, -ct[1]);\n" 250 "float x = fma(ck_rinv[0], x_, fma(ck_rinv[1], y_, ck_rinv[2] * ct1));\n" 251 "float y = fma(ck_rinv[3], x_, fma(ck_rinv[4], y_, ck_rinv[5] * ct1));\n" 252 "float z = fma(ck_rinv[6], x_, fma(ck_rinv[7], y_, ck_rinv[8] * ct1));\n" 253 "if (z != 0)\n" 254 "x /= z, y /= z;\n" 255 "else\n" 256 "x = y = -1;\n" 257 "xmap[0] = x;\n" 258 "ymap[0] = y;\n" 259 "}\n" 260 "}\n" 261 "}\n" 262 "__kernel void buildWarpCylindricalMaps(__global uchar * xmapptr, int xmap_step, int xmap_offset,\n" 263 "__global uchar * ymapptr, int ymap_step, int ymap_offset, int rows, int cols,\n" 264 "__constant float * ck_rinv, int tl_u, int tl_v, float scale, int rowsPerWI)\n" 265 "{\n" 266 "int du = get_global_id(0);\n" 267 "int dv0 = get_global_id(1) * rowsPerWI;\n" 268 "if (du < cols)\n" 269 "{\n" 270 "int xmap_index = mad24(dv0, xmap_step, mad24(du, (int)sizeof(float), xmap_offset));\n" 271 "int ymap_index = mad24(dv0, ymap_step, mad24(du, (int)sizeof(float), ymap_offset));\n" 272 "float u = (tl_u + du) * scale;\n" 273 "float x_, z_;\n" 274 "x_ = sincos(u, &z_);\n" 275 "for (int dv = dv0, dv1 = min(rows, dv0 + rowsPerWI); dv < dv1; ++dv, xmap_index += xmap_step,\n" 276 "ymap_index += ymap_step)\n" 277 "{\n" 278 "__global float * xmap = (__global float *)(xmapptr + xmap_index);\n" 279 "__global float * ymap = (__global float *)(ymapptr + ymap_index);\n" 280 "float y_ = (tl_v + dv) * scale;\n" 281 "float x, y, z;\n" 282 "x = fma(ck_rinv[0], x_, fma(ck_rinv[1], y_, ck_rinv[2] * z_));\n" 283 "y = fma(ck_rinv[3], x_, fma(ck_rinv[4], y_, ck_rinv[5] * z_));\n" 284 "z = fma(ck_rinv[6], x_, fma(ck_rinv[7], y_, ck_rinv[8] * z_));\n" 285 "if (z > 0)\n" 286 "x /= z, y /= z;\n" 287 "else\n" 288 "x = y = -1;\n" 289 "xmap[0] = x;\n" 290 "ymap[0] = y;\n" 291 "}\n" 292 "}\n" 293 "}\n" 294 "__kernel void buildWarpSphericalMaps(__global uchar * xmapptr, int xmap_step, int xmap_offset,\n" 295 "__global uchar * ymapptr, int ymap_step, int ymap_offset, int rows, int cols,\n" 296 "__constant float * ck_rinv, int tl_u, int tl_v, float scale, int rowsPerWI)\n" 297 "{\n" 298 "int du = get_global_id(0);\n" 299 "int dv0 = get_global_id(1) * rowsPerWI;\n" 300 "if (du < cols)\n" 301 "{\n" 302 "int xmap_index = mad24(dv0, xmap_step, mad24(du, (int)sizeof(float), xmap_offset));\n" 303 "int ymap_index = mad24(dv0, ymap_step, mad24(du, (int)sizeof(float), ymap_offset));\n" 304 "float u = (tl_u + du) * scale;\n" 305 "float cosu, sinu = sincos(u, &cosu);\n" 306 "for (int dv = dv0, dv1 = min(rows, dv0 + rowsPerWI); dv < dv1; ++dv, xmap_index += xmap_step,\n" 307 "ymap_index += ymap_step)\n" 308 "{\n" 309 "__global float * xmap = (__global float *)(xmapptr + xmap_index);\n" 310 "__global float * ymap = (__global float *)(ymapptr + ymap_index);\n" 311 "float v = (tl_v + dv) * scale;\n" 312 "float cosv, sinv = sincos(v, &cosv);\n" 313 "float x_ = sinv * sinu;\n" 314 "float y_ = -cosv;\n" 315 "float z_ = sinv * cosu;\n" 316 "float x, y, z;\n" 317 "x = fma(ck_rinv[0], x_, fma(ck_rinv[1], y_, ck_rinv[2] * z_));\n" 318 "y = fma(ck_rinv[3], x_, fma(ck_rinv[4], y_, ck_rinv[5] * z_));\n" 319 "z = fma(ck_rinv[6], x_, fma(ck_rinv[7], y_, ck_rinv[8] * z_));\n" 320 "if (z > 0)\n" 321 "x /= z, y /= z;\n" 322 "else\n" 323 "x = y = -1;\n" 324 "xmap[0] = x;\n" 325 "ymap[0] = y;\n" 326 "}\n" 327 "}\n" 328 "}\n" 329 , "83a61a49d8be5dcc09a00d8d4651c4f8"}; 330 ProgramSource warpers_oclsrc(warpers.programStr); 331 } 332 }} 333