3 "#line 1 \"libavfilter/opencl/convolution.cl\"\n" 5 " * Copyright (c) 2018 Danil Iashchenko\n" 7 " * This file is part of FFmpeg.\n" 9 " * FFmpeg is free software; you can redistribute it and/or\n" 10 " * modify it under the terms of the GNU Lesser General Public\n" 11 " * License as published by the Free Software Foundation; either\n" 12 " * version 2.1 of the License, or (at your option) any later version.\n" 14 " * FFmpeg is distributed in the hope that it will be useful,\n" 15 " * but WITHOUT ANY WARRANTY; without even the implied warranty of\n" 16 " * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU\n" 17 " * Lesser General Public License for more details.\n" 19 " * You should have received a copy of the GNU Lesser General Public\n" 20 " * License along with FFmpeg; if not, write to the Free Software\n" 21 " * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA\n" 24 "__kernel void convolution_global(__write_only image2d_t dst,\n" 25 " __read_only image2d_t src,\n" 26 " int coef_matrix_dim,\n" 27 " __constant float *coef_matrix,\n" 31 " const sampler_t sampler = (CLK_NORMALIZED_COORDS_FALSE |\n" 32 " CLK_ADDRESS_CLAMP_TO_EDGE |\n" 33 " CLK_FILTER_NEAREST);\n" 35 " const int half_matrix_dim = (coef_matrix_dim / 2);\n" 36 " int2 loc = (int2)(get_global_id(0), get_global_id(1));\n" 37 " float4 convPix = (float4)(0.0f, 0.0f, 0.0f, 0.0f);\n" 39 " for (int conv_i = -half_matrix_dim; conv_i <= half_matrix_dim; conv_i++) {\n" 40 " for (int conv_j = -half_matrix_dim; conv_j <= half_matrix_dim; conv_j++) {\n" 41 " float4 px = read_imagef(src, sampler, loc + (int2)(conv_j, conv_i));\n" 42 " convPix += px * coef_matrix[(conv_i + half_matrix_dim) * coef_matrix_dim +\n" 43 " (conv_j + half_matrix_dim)];\n" 46 " float4 dstPix = convPix * div + bias;\n" 47 " write_imagef(dst, loc, dstPix);\n" 51 "__kernel void sobel_global(__write_only image2d_t dst,\n" 52 " __read_only image2d_t src,\n" 56 " const sampler_t sampler = (CLK_NORMALIZED_COORDS_FALSE |\n" 57 " CLK_ADDRESS_CLAMP_TO_EDGE |\n" 58 " CLK_FILTER_NEAREST);\n" 60 " int2 loc = (int2)(get_global_id(0), get_global_id(1));\n" 62 " float4 sum1 = read_imagef(src, sampler, loc + (int2)(-1,-1)) * -1 +\n" 63 " read_imagef(src, sampler, loc + (int2)( 0,-1)) * -2 +\n" 64 " read_imagef(src, sampler, loc + (int2)( 1,-1)) * -1 +\n" 65 " read_imagef(src, sampler, loc + (int2)(-1, 1)) * 1 +\n" 66 " read_imagef(src, sampler, loc + (int2)( 0, 1)) * 2 +\n" 67 " read_imagef(src, sampler, loc + (int2)( 1, 1)) * 1;\n" 69 " float4 sum2 = read_imagef(src, sampler, loc + (int2)(-1,-1)) * -1 +\n" 70 " read_imagef(src, sampler, loc + (int2)(-1, 0)) * -2 +\n" 71 " read_imagef(src, sampler, loc + (int2)(-1, 1)) * -1 +\n" 72 " read_imagef(src, sampler, loc + (int2)( 1,-1)) * 1 +\n" 73 " read_imagef(src, sampler, loc + (int2)( 1, 0)) * 2 +\n" 74 " read_imagef(src, sampler, loc + (int2)( 1, 1)) * 1;\n" 76 " float4 dstPix = hypot(sum1, sum2) * div + bias;\n" 77 " write_imagef(dst, loc, dstPix);\n" 80 "__kernel void prewitt_global(__write_only image2d_t dst,\n" 81 " __read_only image2d_t src,\n" 85 " const sampler_t sampler = (CLK_NORMALIZED_COORDS_FALSE |\n" 86 " CLK_ADDRESS_CLAMP_TO_EDGE |\n" 87 " CLK_FILTER_NEAREST);\n" 89 " int2 loc = (int2)(get_global_id(0), get_global_id(1));\n" 91 " float4 sum1 = read_imagef(src, sampler, loc + (int2)(-1,-1)) * 1 +\n" 92 " read_imagef(src, sampler, loc + (int2)( 0,-1)) * 1 +\n" 93 " read_imagef(src, sampler, loc + (int2)( 1,-1)) * 1 +\n" 94 " read_imagef(src, sampler, loc + (int2)(-1, 1)) * -1 +\n" 95 " read_imagef(src, sampler, loc + (int2)( 0, 1)) * -1 +\n" 96 " read_imagef(src, sampler, loc + (int2)( 1, 1)) * -1;\n" 98 " float4 sum2 = read_imagef(src, sampler, loc + (int2)(-1,-1)) * 1 +\n" 99 " read_imagef(src, sampler, loc + (int2)(-1, 0)) * 1 +\n" 100 " read_imagef(src, sampler, loc + (int2)(-1, 1)) * 1 +\n" 101 " read_imagef(src, sampler, loc + (int2)( 1,-1)) * -1 +\n" 102 " read_imagef(src, sampler, loc + (int2)( 1, 0)) * -1 +\n" 103 " read_imagef(src, sampler, loc + (int2)( 1, 1)) * -1;\n" 105 " float4 dstPix = hypot(sum1, sum2) * div + bias;\n" 106 " write_imagef(dst, loc, dstPix);\n" 109 "__kernel void roberts_global(__write_only image2d_t dst,\n" 110 " __read_only image2d_t src,\n" 114 " const sampler_t sampler = (CLK_NORMALIZED_COORDS_FALSE |\n" 115 " CLK_ADDRESS_CLAMP_TO_EDGE |\n" 116 " CLK_FILTER_NEAREST);\n" 118 " int2 loc = (int2)(get_global_id(0), get_global_id(1));\n" 120 " float4 sum1 = read_imagef(src, sampler, loc + (int2)(-1,-1)) * 1 +\n" 121 " read_imagef(src, sampler, loc + (int2)( 0,-1)) * -1;\n" 124 " float4 sum2 = read_imagef(src, sampler, loc + (int2)(-1, 0)) * -1 +\n" 125 " read_imagef(src, sampler, loc + (int2)( 0, 0)) * 1;\n" 128 " float4 dstPix = hypot(sum1, sum2) * div + bias;\n" 129 " write_imagef(dst, loc, dstPix);\n" const char * ff_opencl_source_convolution