git.sesse.net Git - ffmpeg/blob - libavfilter/opencl/convolution.cl

   1 /*
   2  * Copyright (c) 2018 Danil Iashchenko
   3  *
   4  * This file is part of FFmpeg.
   5  *
   6  * FFmpeg is free software; you can redistribute it and/or
   7  * modify it under the terms of the GNU Lesser General Public
   8  * License as published by the Free Software Foundation; either
   9  * version 2.1 of the License, or (at your option) any later version.
  10  *
  11  * FFmpeg is distributed in the hope that it will be useful,
  12  * but WITHOUT ANY WARRANTY; without even the implied warranty of
  13  * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
  14  * Lesser General Public License for more details.
  15  *
  16  * You should have received a copy of the GNU Lesser General Public
  17  * License along with FFmpeg; if not, write to the Free Software
  18  * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
  19  */
  20
  21 __kernel void convolution_global(__write_only image2d_t dst,
  22                                  __read_only  image2d_t src,
  23                                  int coef_matrix_dim,
  24                                  __constant float *coef_matrix,
  25                                  float div,
  26                                  float bias)
  27 {
  28     const sampler_t sampler = (CLK_NORMALIZED_COORDS_FALSE |
  29                                CLK_ADDRESS_CLAMP_TO_EDGE   |
  30                                CLK_FILTER_NEAREST);
  31
  32     const int half_matrix_dim = (coef_matrix_dim / 2);
  33     int2 loc = (int2)(get_global_id(0), get_global_id(1));
  34     float4 convPix = (float4)(0.0f, 0.0f, 0.0f, 0.0f);
  35
  36     for (int conv_i = -half_matrix_dim; conv_i <= half_matrix_dim; conv_i++) {
  37         for (int conv_j = -half_matrix_dim; conv_j <= half_matrix_dim; conv_j++) {
  38             float4 px = read_imagef(src, sampler, loc + (int2)(conv_j, conv_i));
  39             convPix += px * coef_matrix[(conv_i + half_matrix_dim) * coef_matrix_dim +
  40                                         (conv_j + half_matrix_dim)];
  41         }
  42      }
  43      float4 dstPix = convPix * div + bias;
  44      write_imagef(dst, loc, dstPix);
  45 }