| // This file is auto-generated. Do not edit! |
| |
| #include "precomp.hpp" |
| #include "opencl_kernels_superres.hpp" |
| |
| namespace cv |
| { |
| namespace ocl |
| { |
| namespace superres |
| { |
| |
| const struct ProgramEntry superres_btvl1={"superres_btvl1", |
| "#ifndef cn\n" |
| "#define cn 1\n" |
| "#endif\n" |
| "#define sz (int)sizeof(float)\n" |
| "#define src_elem_at(_src, y, step, x) *(__global const float *)(_src + mad24(y, step, (x) * sz))\n" |
| "#define dst_elem_at(_dst, y, step, x) *(__global float *)(_dst + mad24(y, step, (x) * sz))\n" |
| "__kernel void buildMotionMaps(__global const uchar * forwardMotionPtr, int forwardMotion_step, int forwardMotion_offset,\n" |
| "__global const uchar * backwardMotionPtr, int backwardMotion_step, int backwardMotion_offset,\n" |
| "__global const uchar * forwardMapPtr, int forwardMap_step, int forwardMap_offset,\n" |
| "__global const uchar * backwardMapPtr, int backwardMap_step, int backwardMap_offset,\n" |
| "int rows, int cols)\n" |
| "{\n" |
| "int x = get_global_id(0);\n" |
| "int y = get_global_id(1);\n" |
| "if (x < cols && y < rows)\n" |
| "{\n" |
| "int forwardMotion_index = mad24(forwardMotion_step, y, (int)sizeof(float2) * x + forwardMotion_offset);\n" |
| "int backwardMotion_index = mad24(backwardMotion_step, y, (int)sizeof(float2) * x + backwardMotion_offset);\n" |
| "int forwardMap_index = mad24(forwardMap_step, y, (int)sizeof(float2) * x + forwardMap_offset);\n" |
| "int backwardMap_index = mad24(backwardMap_step, y, (int)sizeof(float2) * x + backwardMap_offset);\n" |
| "float2 forwardMotion = *(__global const float2 *)(forwardMotionPtr + forwardMotion_index);\n" |
| "float2 backwardMotion = *(__global const float2 *)(backwardMotionPtr + backwardMotion_index);\n" |
| "__global float2 * forwardMap = (__global float2 *)(forwardMapPtr + forwardMap_index);\n" |
| "__global float2 * backwardMap = (__global float2 *)(backwardMapPtr + backwardMap_index);\n" |
| "float2 basePoint = (float2)(x, y);\n" |
| "forwardMap[0] = basePoint + backwardMotion;\n" |
| "backwardMap[0] = basePoint + forwardMotion;\n" |
| "}\n" |
| "}\n" |
| "__kernel void upscale(__global const uchar * srcptr, int src_step, int src_offset, int src_rows, int src_cols,\n" |
| "__global uchar * dstptr, int dst_step, int dst_offset, int scale)\n" |
| "{\n" |
| "int x = get_global_id(0);\n" |
| "int y = get_global_id(1);\n" |
| "if (x < src_cols && y < src_rows)\n" |
| "{\n" |
| "int src_index = mad24(y, src_step, sz * x * cn + src_offset);\n" |
| "int dst_index = mad24(y * scale, dst_step, sz * x * scale * cn + dst_offset);\n" |
| "__global const float * src = (__global const float *)(srcptr + src_index);\n" |
| "__global float * dst = (__global float *)(dstptr + dst_index);\n" |
| "#pragma unroll\n" |
| "for (int c = 0; c < cn; ++c)\n" |
| "dst[c] = src[c];\n" |
| "}\n" |
| "}\n" |
| "inline float diffSign1(float a, float b)\n" |
| "{\n" |
| "return a > b ? 1.0f : a < b ? -1.0f : 0.0f;\n" |
| "}\n" |
| "inline float3 diffSign3(float3 a, float3 b)\n" |
| "{\n" |
| "float3 pos;\n" |
| "pos.x = a.x > b.x ? 1.0f : a.x < b.x ? -1.0f : 0.0f;\n" |
| "pos.y = a.y > b.y ? 1.0f : a.y < b.y ? -1.0f : 0.0f;\n" |
| "pos.z = a.z > b.z ? 1.0f : a.z < b.z ? -1.0f : 0.0f;\n" |
| "return pos;\n" |
| "}\n" |
| "__kernel void diffSign(__global const uchar * src1, int src1_step, int src1_offset,\n" |
| "__global const uchar * src2, int src2_step, int src2_offset,\n" |
| "__global uchar * dst, int dst_step, int dst_offset, int rows, int cols)\n" |
| "{\n" |
| "int x = get_global_id(0);\n" |
| "int y = get_global_id(1);\n" |
| "if (x < cols && y < rows)\n" |
| "*(__global float *)(dst + mad24(y, dst_step, sz * x + dst_offset)) =\n" |
| "diffSign1(*(__global const float *)(src1 + mad24(y, src1_step, sz * x + src1_offset)),\n" |
| "*(__global const float *)(src2 + mad24(y, src2_step, sz * x + src2_offset)));\n" |
| "}\n" |
| "__kernel void calcBtvRegularization(__global const uchar * src, int src_step, int src_offset,\n" |
| "__global uchar * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols,\n" |
| "int ksize, __constant float * c_btvRegWeights)\n" |
| "{\n" |
| "int x = get_global_id(0) + ksize;\n" |
| "int y = get_global_id(1) + ksize;\n" |
| "if (y < dst_rows - ksize && x < dst_cols - ksize)\n" |
| "{\n" |
| "src += src_offset;\n" |
| "#if cn == 1\n" |
| "const float srcVal = src_elem_at(src, y, src_step, x);\n" |
| "float dstVal = 0.0f;\n" |
| "for (int m = 0, count = 0; m <= ksize; ++m)\n" |
| "for (int l = ksize; l + m >= 0; --l, ++count)\n" |
| "{\n" |
| "dstVal += c_btvRegWeights[count] * (diffSign1(srcVal, src_elem_at(src, y + m, src_step, x + l))\n" |
| "- diffSign1(src_elem_at(src, y - m, src_step, x - l), srcVal));\n" |
| "}\n" |
| "dst_elem_at(dst, y, dst_step, x) = dstVal;\n" |
| "#elif cn == 3\n" |
| "__global const float * src0ptr = (__global const float *)(src + mad24(y, src_step, 3 * sz * x + src_offset));\n" |
| "float3 srcVal = (float3)(src0ptr[0], src0ptr[1], src0ptr[2]), dstVal = 0.f;\n" |
| "for (int m = 0, count = 0; m <= ksize; ++m)\n" |
| "{\n" |
| "for (int l = ksize; l + m >= 0; --l, ++count)\n" |
| "{\n" |
| "__global const float * src1ptr = (__global const float *)(src + mad24(y + m, src_step, 3 * sz * (x + l) + src_offset));\n" |
| "__global const float * src2ptr = (__global const float *)(src + mad24(y - m, src_step, 3 * sz * (x - l) + src_offset));\n" |
| "float3 src1 = (float3)(src1ptr[0], src1ptr[1], src1ptr[2]);\n" |
| "float3 src2 = (float3)(src2ptr[0], src2ptr[1], src2ptr[2]);\n" |
| "dstVal += c_btvRegWeights[count] * (diffSign3(srcVal, src1) - diffSign3(src2, srcVal));\n" |
| "}\n" |
| "}\n" |
| "__global float * dstptr = (__global float *)(dst + mad24(y, dst_step, 3 * sz * x + dst_offset + 0));\n" |
| "dstptr[0] = dstVal.x;\n" |
| "dstptr[1] = dstVal.y;\n" |
| "dstptr[2] = dstVal.z;\n" |
| "#else\n" |
| "#error \"Number of channels should be either 1 of 3\"\n" |
| "#endif\n" |
| "}\n" |
| "}\n" |
| , "747d707919dd581b14986fa59ef19f88"}; |
| ProgramSource superres_btvl1_oclsrc(superres_btvl1.programStr); |
| } |
| }} |
