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  */
  34
  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
  41
  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 }
  51
  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 }
  65
  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 }
  72
  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 }
  91
  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;
 102
 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     }
 108
 109   /* process the interior pixels */
 110   for (k = 2; k < height; k++) {
 111     dest += dstride;
 112     src += sstride;
 113
 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   }
 135 }
 136 /* nothing past this point */