Add OpenCL video kernels for median and sobel img processing functions. They are...
[processor-sdk/gst-plugin-dsp66.git] / src / kernels / oclconv / conv.cl
1 /******************************************************************************
2  * Copyright (c) 2013-2014, Texas Instruments Incorporated - http://www.ti.com/
3  *   All rights reserved.
4  *
5  *   Redistribution and use in source and binary forms, with or without
6  *   modification, are permitted provided that the following conditions are met:
7  *       * Redistributions of source code must retain the above copyright
8  *         notice, this list of conditions and the following disclaimer.
9  *       * Redistributions in binary form must reproduce the above copyright
10  *         notice, this list of conditions and the following disclaimer in the
11  *         documentation and/or other materials provided with the distribution.
12  *       * Neither the name of Texas Instruments Incorporated nor the
13  *         names of its contributors may be used to endorse or promote products
14  *         derived from this software without specific prior written permission.
15  *
16  *   THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
17  *   AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
18  *   IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE 
19  *   ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR CONTRIBUTORS BE
20  *   LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
21  *   CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
22  *   SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
23  *   INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
24  *   CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
25  *   ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
26  *   THE POSSIBILITY OF SUCH DAMAGE.
27  *****************************************************************************/
28 #define PIX_SORT(a,b) { if ((a)>(b)) PIX_SWAP((a),(b)); }
29 #define PIX_SWAP(a,b) { unsigned char temp=(a);(a)=(b);(b)=temp; }
31 void IMG_median_3x3_8 (const unsigned char *restrict in_data, int cols, unsigned char * restrict out_data);
32 void IMG_sobel_3x3_8  (const unsigned char *restrict in_data, unsigned char * restrict out_data, int rows, int cols);
34 kernel void Median3x3(global const uchar* src, global uchar *dest,
35                       const int width, const int height,
36                       const int dstride, const int sstride)
37 {
38   int i;
39   const int max_X = width - 1;
40   const int max_Y = height - 1;
41     for (i = 0; i < max_Y; i++) {
42       IMG_median_3x3_8 ((const unsigned char *)src, max_X, (const unsigned char *)dest);
43       src  += sstride;
44       dest += dstride; 
45     }
46 }
48 kernel void Sobel3x3(global const uchar* src, global uchar *dest,
49                       const int width, const int height,
50                       const int dstride, const int sstride)
51 {
52   IMG_sobel_3x3_8((const unsigned char *)src, (const unsigned char *)dest, width, height);
53 }
55 kernel void Median2x2 (global const uchar* src, global uchar *dest,
56                       const int width, const int height,
57                       const int dstride, const int sstride)
58 {
59 /***
60     int id = get_global_id(0);
61     c[id] = a[id] + b[id];
62 ***/
63   unsigned char p[5];
64   int i, j, k;
66     /* copy the top and bottom rows into the result array */
67     for (i = 0; i < width; i++) {
68       dest[i] = src[i];
69       dest[(height - 1) * dstride + i] = src[(height - 1) * sstride + i];
70     }
72   /* process the interior pixels */
73   for (k = 2; k < height; k++) {
74     dest += dstride;
75     src += sstride;
77     dest[0] = src[0];
78     for (j = 2, i = 1; j < width; j++, i++) {
79       p[0] = src[i - sstride];
80       p[1] = src[i - 1];
81       p[2] = src[i];
82       p[3] = src[i + 1];
83       p[4] = src[i + sstride];
84       PIX_SORT (p[0], p[1]);
85       PIX_SORT (p[3], p[4]);
86       PIX_SORT (p[0], p[3]);
87       PIX_SORT (p[1], p[4]);
88       PIX_SORT (p[1], p[2]);
89       PIX_SORT (p[2], p[3]);
90       PIX_SORT (p[1], p[2]);
91       /* debug line */
92       if(k > 24 && k < 32) dest[i] = 0;
93       else
94       dest[i] = p[2];
95     }
96     dest[i] = src[i];
97   }
98 }
99 /* nothing past this point */