modules/imgproc/src/opencl/gaussianBlur3x3.cl

   1 // This file is part of OpenCV project.
   2 // It is subject to the license terms in the LICENSE file found in the top-level directory
   3 // of this distribution and at http://opencv.org/license.html.
   4
   5 #define DIG(a) a,
   6 __constant float kx[] = { KERNEL_MATRIX_X };
   7 __constant float ky[] = { KERNEL_MATRIX_Y };
   8
   9 #define OP(delta, y, x) (convert_float16(arr[(y + delta) * 3 + x]) * ky[y] * kx[x])
  10
  11 __kernel void gaussianBlur3x3_8UC1_cols16_rows2(__global const uint* src, int src_step,
  12                                                 __global uint* dst, int dst_step, int rows, int cols)
  13 {
  14     int block_x = get_global_id(0);
  15     int y = get_global_id(1) * 2;
  16     int ssx, dsx;
  17
  18     if ((block_x * 16) >= cols || y >= rows) return;
  19
  20     uint4 line[4];
  21     uint4 line_out[2];
  22     uchar a; uchar16 b; uchar c;
  23     uchar d; uchar16 e; uchar f;
  24     uchar g; uchar16 h; uchar i;
  25     uchar j; uchar16 k; uchar l;
  26
  27     ssx = dsx = 1;
  28     int src_index = block_x * 4 * ssx + (y - 1) * (src_step / 4);
  29     line[1] = vload4(0, src + src_index + (src_step / 4));
  30     line[2] = vload4(0, src + src_index + 2 * (src_step / 4));
  31
  32 #ifdef BORDER_CONSTANT
  33     line[0] = (y == 0) ? (uint4)0 : vload4(0, src + src_index);
  34     line[3] = (y == (rows - 2)) ? (uint4)0 : vload4(0, src + src_index + 3 * (src_step / 4));
  35 #elif defined BORDER_REFLECT_101
  36     line[0] = (y == 0) ? line[2] : vload4(0, src + src_index);
  37     line[3] = (y == (rows - 2)) ? line[1] : vload4(0, src + src_index + 3 * (src_step / 4));
  38 #elif defined (BORDER_REPLICATE) || defined(BORDER_REFLECT)
  39     line[0] = (y == 0) ? line[1] : vload4(0, src + src_index);
  40     line[3] = (y == (rows - 2)) ? line[2] : vload4(0, src + src_index + 3 * (src_step / 4));
  41 #endif
  42
  43     __global uchar *src_p = (__global uchar *)src;
  44
  45     src_index = block_x * 16 * ssx + (y - 1) * src_step;
  46     bool line_end = ((block_x + 1) * 16 == cols);
  47
  48     b = as_uchar16(line[0]);
  49     e = as_uchar16(line[1]);
  50     h = as_uchar16(line[2]);
  51     k = as_uchar16(line[3]);
  52
  53 #ifdef BORDER_CONSTANT
  54     a = (block_x == 0 || y == 0) ? 0 : src_p[src_index - 1];
  55     c = (line_end || y == 0) ? 0 : src_p[src_index + 16];
  56
  57     d = (block_x == 0) ? 0 : src_p[src_index + src_step - 1];
  58     f = line_end ? 0 : src_p[src_index + src_step + 16];
  59
  60     g = (block_x == 0) ? 0 : src_p[src_index + 2 * src_step - 1];
  61     i = line_end ? 0 : src_p[src_index + 2 * src_step + 16];
  62
  63     j = (block_x == 0 || y == (rows - 2)) ? 0 : src_p[src_index + 3 * src_step - 1];
  64     l = (line_end || y == (rows - 2))? 0 : src_p[src_index + 3 * src_step + 16];
  65
  66 #elif defined BORDER_REFLECT_101
  67     int offset;
  68     offset = (y == 0) ? (2 * src_step) : 0;
  69
  70     a = (block_x == 0) ? src_p[src_index + offset + 1] : src_p[src_index + offset - 1];
  71     c = line_end ? src_p[src_index + offset + 14] : src_p[src_index + offset + 16];
  72
  73     d = (block_x == 0) ? src_p[src_index + src_step + 1] : src_p[src_index + src_step - 1];
  74     f = line_end ? src_p[src_index + src_step + 14] : src_p[src_index + src_step + 16];
  75
  76     g = (block_x == 0) ? src_p[src_index + 2 * src_step + 1] : src_p[src_index + 2 * src_step - 1];
  77     i = line_end ? src_p[src_index + 2 * src_step + 14] : src_p[src_index + 2 * src_step + 16];
  78
  79     offset = (y == (rows - 2)) ? (1 * src_step) : (3 * src_step);
  80
  81     j = (block_x == 0) ? src_p[src_index + offset + 1] : src_p[src_index + offset - 1];
  82     l = line_end ? src_p[src_index + offset + 14] : src_p[src_index + offset + 16];
  83
  84 #elif defined (BORDER_REPLICATE) || defined(BORDER_REFLECT)
  85     int offset;
  86     offset = (y == 0) ? (1 * src_step) : 0;
  87
  88     a = (block_x == 0) ? src_p[src_index + offset] : src_p[src_index + offset - 1];
  89     c = line_end ? src_p[src_index + offset + 15] : src_p[src_index + offset + 16];
  90
  91     d = (block_x == 0) ? src_p[src_index + src_step] : src_p[src_index + src_step - 1];
  92     f = line_end ? src_p[src_index + src_step + 15] : src_p[src_index + src_step + 16];
  93
  94     g = (block_x == 0) ? src_p[src_index + 2 * src_step] : src_p[src_index + 2 * src_step - 1];
  95     i = line_end ? src_p[src_index + 2 * src_step + 15] : src_p[src_index + 2 * src_step + 16];
  96
  97     offset = (y == (rows - 2)) ? (2 * src_step) : (3 * src_step);
  98
  99     j = (block_x == 0) ? src_p[src_index + offset] : src_p[src_index + offset - 1];
 100     l = line_end ? src_p[src_index + offset + 15] : src_p[src_index + offset + 16];
 101 #endif
 102
 103     uchar16 arr[12];
 104     float16 sum[2];
 105
 106     arr[0] = (uchar16)(a, b.s0123, b.s456789ab, b.scde);
 107     arr[1] = b;
 108     arr[2] = (uchar16)(b.s123, b.s4567, b.s89abcdef, c);
 109     arr[3] = (uchar16)(d, e.s0123, e.s456789ab, e.scde);
 110     arr[4] = e;
 111     arr[5] = (uchar16)(e.s123, e.s4567, e.s89abcdef, f);
 112     arr[6] = (uchar16)(g, h.s0123, h.s456789ab, h.scde);
 113     arr[7] = h;
 114     arr[8] = (uchar16)(h.s123, h.s4567, h.s89abcdef, i);
 115     arr[9] = (uchar16)(j, k.s0123, k.s456789ab, k.scde);
 116     arr[10] = k;
 117     arr[11] = (uchar16)(k.s123, k.s4567, k.s89abcdef, l);
 118
 119     sum[0] = OP(0, 0, 0) + OP(0, 0, 1) + OP(0, 0, 2) +
 120              OP(0, 1, 0) + OP(0, 1, 1) + OP(0, 1, 2) +
 121              OP(0, 2, 0) + OP(0, 2, 1) + OP(0, 2, 2);
 122
 123     sum[1] = OP(1, 0, 0) + OP(1, 0, 1) + OP(1, 0, 2) +
 124              OP(1, 1, 0) + OP(1, 1, 1) + OP(1, 1, 2) +
 125              OP(1, 2, 0) + OP(1, 2, 1) + OP(1, 2, 2);
 126
 127     line_out[0] = as_uint4(convert_uchar16_sat_rte(sum[0]));
 128     line_out[1] = as_uint4(convert_uchar16_sat_rte(sum[1]));
 129
 130     int dst_index = block_x * 4 * dsx + y * (dst_step / 4);
 131     vstore4(line_out[0], 0, dst + dst_index);
 132     vstore4(line_out[1], 0, dst + dst_index + (dst_step / 4));
 133 }