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