Copyright cleanup - fix so we have BSD licensing banners
[processor-sdk/gst-plugin-dsp66.git] / src / kernels / oclconv / conv.cl
1 /*
2  * Copyright (C) 2014 Texas Instruments Incorporated - http://www.ti.com/
3  *
4  *
5  *  Redistribution and use in source and binary forms, with or without
6  *  modification, are permitted provided that the following conditions
7  *  are met:
8  *
9  *    Redistributions of source code must retain the above copyright
10  *    notice, this list of conditions and the following disclaimer.
11  *
12  *    Redistributions in binary form must reproduce the above copyright
13  *    notice, this list of conditions and the following disclaimer in the
14  *    documentation and/or other materials provided with the  
15  *    distribution.
16  *
17  *    Neither the name of Texas Instruments Incorporated nor the names of
18  *    its contributors may be used to endorse or promote products derived
19  *    from this software without specific prior written permission.
20  *
21  *  THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
22  *  "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT
23  *  LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR
24  *  A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT
25  *  OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL,
26  *  SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT
27  *  LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE,
28  *  DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY
29  *  THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
30  *  (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
31  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
32  *
33  */
35 #define PIX_SORT(a,b) { if ((a)>(b)) PIX_SWAP((a),(b)); }
36 #define PIX_SWAP(a,b) { unsigned char temp=(a);(a)=(b);(b)=temp; }
37 #define ARGS global uchar  *pInput,    global ushort *pBufGradX, \
38              global ushort *pBufGradY, global ushort *pBufMag,   \
39              global uchar  *pBufOut,   global uchar  *pScratch,  \
40              global uint   *numItems,  ushort width, ushort height
42 /* imglib library functions */
43 void IMG_median_3x3_8 (const unsigned char *restrict in_data, int cols, unsigned char * restrict out_data);
44 void IMG_sobel_3x3_8  (const unsigned char *restrict in_data, unsigned char *restrict out_data, int rows, int cols);
45 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);
46 void VLIB_Canny_Edge_Detection (ARGS);
47 kernel void canny_tiocl(ARGS)
48 {
49     VLIB_Canny_Edge_Detection(pInput, pBufGradX, pBufGradY, pBufMag, pBufOut, pScratch, numItems, width, height);
50 }
52 kernel void Median3x3(global const uchar* src, global uchar *dest,
53                       const int width, const int height,
54                       const int dstride, const int sstride)
55 {
56   int i;
57   const int max_X = width - 1;
58   const int max_Y = height - 1;
59     for (i = 0; i < max_Y; i++) {
60       IMG_median_3x3_8 ((const unsigned char *)src, max_X, (const unsigned char *)dest);
61       src  += sstride;
62       dest += dstride; 
63     }
64 }
66 kernel void Sobel3x3(global const uchar* src, global uchar *dest,
67                       const int width, const int height,
68                       const int dstride, const int sstride)
69 {
70   IMG_sobel_3x3_8((const unsigned char *)src, (const unsigned char *)dest, width, height);
71 }
73 kernel void Conv5x5(global const uchar* src, global uchar *dest,
74                       const int width, const int height,
75                       const int dstride, const int sstride)
76 {
77   int i;
78   const char conv_mask5x5[25] = {
79      1, 4, 6, 4, 1,
80      4,16,24,16, 4,
81      6,24,36,24, 6,
82      4,16,24,16, 4,
83      1, 4, 6, 4, 1
84   };
85     for (i = 0; i < height; i++) {
86       IMG_conv_5x5_i8_c8s ((const unsigned char *)src, (unsigned char *)dest, width, sstride, conv_mask5x5, 8); 
87       src  += sstride;
88       dest += dstride; 
89     }
90 }
92 kernel void Median2x2 (global const uchar* src, global uchar *dest,
93                       const int width, const int height,
94                       const int dstride, const int sstride)
95 {
96 /***
97     int id = get_global_id(0);
98     c[id] = a[id] + b[id];
99 ***/
100   unsigned char p[5];
101   int i, j, k;
103     /* copy the top and bottom rows into the result array */
104     for (i = 0; i < width; i++) {
105       dest[i] = src[i];
106       dest[(height - 1) * dstride + i] = src[(height - 1) * sstride + i];
107     }
109   /* process the interior pixels */
110   for (k = 2; k < height; k++) {
111     dest += dstride;
112     src += sstride;
114     dest[0] = src[0];
115     for (j = 2, i = 1; j < width; j++, i++) {
116       p[0] = src[i - sstride];
117       p[1] = src[i - 1];
118       p[2] = src[i];
119       p[3] = src[i + 1];
120       p[4] = src[i + sstride];
121       PIX_SORT (p[0], p[1]);
122       PIX_SORT (p[3], p[4]);
123       PIX_SORT (p[0], p[3]);
124       PIX_SORT (p[1], p[4]);
125       PIX_SORT (p[1], p[2]);
126       PIX_SORT (p[2], p[3]);
127       PIX_SORT (p[1], p[2]);
128       /* debug line */
129       if(k > 24 && k < 32) dest[i] = 0;
130       else
131       dest[i] = p[2];
132     }
133     dest[i] = src[i];
134   }
136 /* nothing past this point */