815b5ff0e21af96c9d418bfbc14f596b46a0f3fe
[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; }
30 #define ARGS global uchar  *pInput,    global ushort *pBufGradX, \
31              global ushort *pBufGradY, global ushort *pBufMag,   \
32              global uchar  *pBufOut,   global uchar  *pScratch,  \
33              global uint   *numItems,  ushort width, ushort height
35 /* imglib library functions */
36 void IMG_median_3x3_8 (const unsigned char *restrict in_data, int cols, unsigned char * restrict out_data);
37 void IMG_sobel_3x3_8  (const unsigned char *restrict in_data, unsigned char *restrict out_data, int rows, int cols);
38 void IMG_conv_3x3_i8_c8s (const unsigned char *restrict in_data, unsigned char *restrict out_data, int cols, const char *restrict mask, int shift);
39 void VLIB_Canny_Edge_Detection (ARGS);
40 kernel void canny_tiocl(ARGS)
41 {
42     VLIB_Canny_Edge_Detection(pInput, pBufGradX, pBufGradY, pBufMag, pBufOut, pScratch, numItems, width, height);
43 }
45 kernel void Median3x3(global const uchar* src, global uchar *dest,
46                       const int width, const int height,
47                       const int dstride, const int sstride)
48 {
49   int i;
50   const int max_X = width - 1;
51   const int max_Y = height - 1;
52     for (i = 0; i < max_Y; i++) {
53       IMG_median_3x3_8 ((const unsigned char *)src, max_X, (const unsigned char *)dest);
54       src  += sstride;
55       dest += dstride; 
56     }
57 }
59 kernel void Sobel3x3(global const uchar* src, global uchar *dest,
60                       const int width, const int height,
61                       const int dstride, const int sstride)
62 {
63   IMG_sobel_3x3_8((const unsigned char *)src, (const unsigned char *)dest, width, height);
64 }
66 kernel void Conv5x5(global const uchar* src, global uchar *dest,
67                       const int width, const int height,
68                       const int dstride, const int sstride)
69 {
70   int i;
71   const char conv_mask5x5[25] = {
72      1, 4, 6, 4, 1,
73      4,16,24,16, 4,
74      6,24,36,24, 6,
75      4,16,24,16, 4,
76      1, 4, 6, 4, 1
77   };
78     for (i = 0; i < height; i++) {
79       IMG_conv_5x5_i8_c8s ((const unsigned char *)src, (unsigned char *)dest, width, sstride, conv_mask5x5, 8); 
80       src  += sstride;
81       dest += dstride; 
82     }
83 }
85 kernel void Median2x2 (global const uchar* src, global uchar *dest,
86                       const int width, const int height,
87                       const int dstride, const int sstride)
88 {
89 /***
90     int id = get_global_id(0);
91     c[id] = a[id] + b[id];
92 ***/
93   unsigned char p[5];
94   int i, j, k;
96     /* copy the top and bottom rows into the result array */
97     for (i = 0; i < width; i++) {
98       dest[i] = src[i];
99       dest[(height - 1) * dstride + i] = src[(height - 1) * sstride + i];
100     }
102   /* process the interior pixels */
103   for (k = 2; k < height; k++) {
104     dest += dstride;
105     src += sstride;
107     dest[0] = src[0];
108     for (j = 2, i = 1; j < width; j++, i++) {
109       p[0] = src[i - sstride];
110       p[1] = src[i - 1];
111       p[2] = src[i];
112       p[3] = src[i + 1];
113       p[4] = src[i + sstride];
114       PIX_SORT (p[0], p[1]);
115       PIX_SORT (p[3], p[4]);
116       PIX_SORT (p[0], p[3]);
117       PIX_SORT (p[1], p[4]);
118       PIX_SORT (p[1], p[2]);
119       PIX_SORT (p[2], p[3]);
120       PIX_SORT (p[1], p[2]);
121       /* debug line */
122       if(k > 24 && k < 32) dest[i] = 0;
123       else
124       dest[i] = p[2];
125     }
126     dest[i] = src[i];
127   }
129 /* nothing past this point */