1 // This file is auto-generated. Do not edit!
2 
3 #include "precomp.hpp"
4 #include "opencl_kernels_superres.hpp"
5 
6 namespace cv
7 {
8 namespace ocl
9 {
10 namespace superres
11 {
12 
13 const struct ProgramEntry superres_btvl1={"superres_btvl1",
14 "#ifndef cn\n"
15 "#define cn 1\n"
16 "#endif\n"
17 "#define sz (int)sizeof(float)\n"
18 "#define src_elem_at(_src, y, step, x) *(__global const float *)(_src + mad24(y, step, (x) * sz))\n"
19 "#define dst_elem_at(_dst, y, step, x) *(__global float *)(_dst + mad24(y, step, (x) * sz))\n"
20 "__kernel void buildMotionMaps(__global const uchar * forwardMotionPtr, int forwardMotion_step, int forwardMotion_offset,\n"
21 "__global const uchar * backwardMotionPtr, int backwardMotion_step, int backwardMotion_offset,\n"
22 "__global const uchar * forwardMapPtr, int forwardMap_step, int forwardMap_offset,\n"
23 "__global const uchar * backwardMapPtr, int backwardMap_step, int backwardMap_offset,\n"
24 "int rows, int cols)\n"
25 "{\n"
26 "int x = get_global_id(0);\n"
27 "int y = get_global_id(1);\n"
28 "if (x < cols && y < rows)\n"
29 "{\n"
30 "int forwardMotion_index = mad24(forwardMotion_step, y, (int)sizeof(float2) * x + forwardMotion_offset);\n"
31 "int backwardMotion_index = mad24(backwardMotion_step, y, (int)sizeof(float2) * x + backwardMotion_offset);\n"
32 "int forwardMap_index = mad24(forwardMap_step, y, (int)sizeof(float2) * x + forwardMap_offset);\n"
33 "int backwardMap_index = mad24(backwardMap_step, y, (int)sizeof(float2) * x + backwardMap_offset);\n"
34 "float2 forwardMotion = *(__global const float2 *)(forwardMotionPtr + forwardMotion_index);\n"
35 "float2 backwardMotion = *(__global const float2 *)(backwardMotionPtr + backwardMotion_index);\n"
36 "__global float2 * forwardMap = (__global float2 *)(forwardMapPtr + forwardMap_index);\n"
37 "__global float2 * backwardMap = (__global float2 *)(backwardMapPtr + backwardMap_index);\n"
38 "float2 basePoint = (float2)(x, y);\n"
39 "forwardMap[0] = basePoint + backwardMotion;\n"
40 "backwardMap[0] = basePoint + forwardMotion;\n"
41 "}\n"
42 "}\n"
43 "__kernel void upscale(__global const uchar * srcptr, int src_step, int src_offset, int src_rows, int src_cols,\n"
44 "__global uchar * dstptr, int dst_step, int dst_offset, int scale)\n"
45 "{\n"
46 "int x = get_global_id(0);\n"
47 "int y = get_global_id(1);\n"
48 "if (x < src_cols && y < src_rows)\n"
49 "{\n"
50 "int src_index = mad24(y, src_step, sz * x * cn + src_offset);\n"
51 "int dst_index = mad24(y * scale, dst_step, sz * x * scale * cn + dst_offset);\n"
52 "__global const float * src = (__global const float *)(srcptr + src_index);\n"
53 "__global float * dst = (__global float *)(dstptr + dst_index);\n"
54 "#pragma unroll\n"
55 "for (int c = 0; c < cn; ++c)\n"
56 "dst[c] = src[c];\n"
57 "}\n"
58 "}\n"
59 "inline float diffSign1(float a, float b)\n"
60 "{\n"
61 "return a > b ? 1.0f : a < b ? -1.0f : 0.0f;\n"
62 "}\n"
63 "inline float3 diffSign3(float3 a, float3 b)\n"
64 "{\n"
65 "float3 pos;\n"
66 "pos.x = a.x > b.x ? 1.0f : a.x < b.x ? -1.0f : 0.0f;\n"
67 "pos.y = a.y > b.y ? 1.0f : a.y < b.y ? -1.0f : 0.0f;\n"
68 "pos.z = a.z > b.z ? 1.0f : a.z < b.z ? -1.0f : 0.0f;\n"
69 "return pos;\n"
70 "}\n"
71 "__kernel void diffSign(__global const uchar * src1, int src1_step, int src1_offset,\n"
72 "__global const uchar * src2, int src2_step, int src2_offset,\n"
73 "__global uchar * dst, int dst_step, int dst_offset, int rows, int cols)\n"
74 "{\n"
75 "int x = get_global_id(0);\n"
76 "int y = get_global_id(1);\n"
77 "if (x < cols && y < rows)\n"
78 "*(__global float *)(dst + mad24(y, dst_step, sz * x + dst_offset)) =\n"
79 "diffSign1(*(__global const float *)(src1 + mad24(y, src1_step, sz * x + src1_offset)),\n"
80 "*(__global const float *)(src2 + mad24(y, src2_step, sz * x + src2_offset)));\n"
81 "}\n"
82 "__kernel void calcBtvRegularization(__global const uchar * src, int src_step, int src_offset,\n"
83 "__global uchar * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols,\n"
84 "int ksize, __constant float * c_btvRegWeights)\n"
85 "{\n"
86 "int x = get_global_id(0) + ksize;\n"
87 "int y = get_global_id(1) + ksize;\n"
88 "if (y < dst_rows - ksize && x < dst_cols - ksize)\n"
89 "{\n"
90 "src += src_offset;\n"
91 "#if cn == 1\n"
92 "const float srcVal = src_elem_at(src, y, src_step, x);\n"
93 "float dstVal = 0.0f;\n"
94 "for (int m = 0, count = 0; m <= ksize; ++m)\n"
95 "for (int l = ksize; l + m >= 0; --l, ++count)\n"
96 "{\n"
97 "dstVal += c_btvRegWeights[count] * (diffSign1(srcVal, src_elem_at(src, y + m, src_step, x + l))\n"
98 "- diffSign1(src_elem_at(src, y - m, src_step, x - l), srcVal));\n"
99 "}\n"
100 "dst_elem_at(dst, y, dst_step, x) = dstVal;\n"
101 "#elif cn == 3\n"
102 "__global const float * src0ptr = (__global const float *)(src + mad24(y, src_step, 3 * sz * x + src_offset));\n"
103 "float3 srcVal = (float3)(src0ptr[0], src0ptr[1], src0ptr[2]), dstVal = 0.f;\n"
104 "for (int m = 0, count = 0; m <= ksize; ++m)\n"
105 "{\n"
106 "for (int l = ksize; l + m >= 0; --l, ++count)\n"
107 "{\n"
108 "__global const float * src1ptr = (__global const float *)(src + mad24(y + m, src_step, 3 * sz * (x + l) + src_offset));\n"
109 "__global const float * src2ptr = (__global const float *)(src + mad24(y - m, src_step, 3 * sz * (x - l) + src_offset));\n"
110 "float3 src1 = (float3)(src1ptr[0], src1ptr[1], src1ptr[2]);\n"
111 "float3 src2 = (float3)(src2ptr[0], src2ptr[1], src2ptr[2]);\n"
112 "dstVal += c_btvRegWeights[count] * (diffSign3(srcVal, src1) - diffSign3(src2, srcVal));\n"
113 "}\n"
114 "}\n"
115 "__global float * dstptr = (__global float *)(dst + mad24(y, dst_step, 3 * sz * x + dst_offset + 0));\n"
116 "dstptr[0] = dstVal.x;\n"
117 "dstptr[1] = dstVal.y;\n"
118 "dstptr[2] = dstVal.z;\n"
119 "#else\n"
120 "#error \"Number of channels should be either 1 of 3\"\n"
121 "#endif\n"
122 "}\n"
123 "}\n"
124 , "747d707919dd581b14986fa59ef19f88"};
125 ProgramSource superres_btvl1_oclsrc(superres_btvl1.programStr);
126 }
127 }}
128