Add OpenCL video kernels for median and sobel img processing functions. They are...

diff --git a/src/kernels/make.inc b/src/kernels/make.inc
new file mode 100644 (file)
index 0000000..797e64b
--- /dev/null
+++ b/src/kernels/make.inc
@@ -0,0 +1,79 @@
+DSP_INCLUDE  = -I$(TI_OCL_CGT_INSTALL)/include 
+DSP_INCLUDE += -I$(TARGET_ROOTDIR)/usr/share/ti/cgt-c6x/include
+DSP_INCLUDE += -I$(TARGET_ROOTDIR)/usr/share/ti/opencl
+TI_IMGLIB_DIR=${TARGET_ROOTDIR}/usr/share/ti/ti-imglib-c66x-tree
+TI_VLIB_DIR=${TARGET_ROOTDIR}/usr/share/ti/ti-vlib-c66x-tree
+
+CPP   = g++ 
+CL6X  = cl6x -mv6600 --abi=eabi $(DSP_INCLUDE)
+CLOCL = clocl
+
+LIBS  = -lOpenCL -locl_util
+
+UNAME_M :=$(shell uname -m)
+
+# ----------------------------------------------------------------------------
+# If TI_OCL_INSTALL is set, setup make from that location
+# ----------------------------------------------------------------------------
+ifneq ($(TI_OCL_INSTALL),)
+    CPP        = g++
+    CPP_FLAGS += -I$(TI_OCL_INSTALL)/usr/include
+    LD_FLAGS  += -L$(TI_OCL_INSTALL)/usr/lib 
+    LIBS      += -lbfd
+
+# ----------------------------------------------------------------------------
+# Otherwise, if making on x86 assume cross compile for Arm host
+# ----------------------------------------------------------------------------
+else ifneq (,$(findstring 86, $(UNAME_M)))
+    .DEFAULT_GOAL := cross
+
+    # In a cross compile environment we are assuming that the EVM file system
+    # is located on the build host and necessary ARM libraries are installed
+    # on that file system. 
+    ifneq ($(MAKECMDGOALS),clean)
+       ifeq ($(TARGET_ROOTDIR),)
+         $(error Environment variable TARGET_ROOTDIR must be defined. Set it to point at the EVM root file system)
+       endif
+    endif
+
+    # gcc ARM cross compiler will not, by default, search the host's
+    # /usr/include.  Explicitly specify here to find dependent vendor headers
+    cross: override CPP = arm-linux-gnueabihf-g++ 
+    cross: CPP_FLAGS += -I$(TARGET_ROOTDIR)/usr/include -idirafter /usr/include
+
+    # If cross-compilineg, provide path to dependent ARM libraries on the 
+    # target filesystem
+    cross: LD_FLAGS = -L$(TARGET_ROOTDIR)/lib -L$(TARGET_ROOTDIR)/usr/lib -Wl,-rpath-link,$(TARGET_ROOTDIR)/lib -Wl,-rpath-link,$(TARGET_ROOTDIR)/usr/lib 
+endif
+
+%.o: %.cpp
+       @$(CPP) -c $(CPP_FLAGS) $<
+       @echo Compiling $<
+
+%.o: %.c
+       @$(CPP) -c $(CPP_FLAGS) $<
+       @echo Compiling $<
+
+%.obj: %.c
+       @$(CL6X) -c $(CL6X_FLAGS) $<
+       @echo Compiling $<
+
+%.out: %.cl
+       @$(CLOCL) $(CLOCL_FLAGS) $^
+       @echo Compiling $< 
+
+%.dsp_h: %.cl
+       @$(CLOCL) -t $(CLOCL_FLAGS) $^
+       @echo Compiling $< 
+
+$(EXE):
+
+cross:     $(EXE)
+
+clean::
+       @rm -f $(EXE) *.o *.obj *.out *.asm *.if *.opt *.bc *.objc *.map *.bin *.dsp_h
+
+test: clean $(EXE)
+       @echo Running   $(EXE)
+       @./$(EXE) >> /dev/null
+       @if [ $$? -ne 0 ] ; then echo "FAILED !!!" ; fi

diff --git a/src/kernels/oclconv/conv.cl b/src/kernels/oclconv/conv.cl
new file mode 100644 (file)
index 0000000..bc9aced
--- /dev/null
+++ b/src/kernels/oclconv/conv.cl
@@ -0,0 +1,99 @@
+/******************************************************************************
+ * Copyright (c) 2013-2014, Texas Instruments Incorporated - http://www.ti.com/
+ *   All rights reserved.
+ *
+ *   Redistribution and use in source and binary forms, with or without
+ *   modification, are permitted provided that the following conditions are met:
+ *       * Redistributions of source code must retain the above copyright
+ *         notice, this list of conditions and the following disclaimer.
+ *       * Redistributions in binary form must reproduce the above copyright
+ *         notice, this list of conditions and the following disclaimer in the
+ *         documentation and/or other materials provided with the distribution.
+ *       * Neither the name of Texas Instruments Incorporated nor the
+ *         names of its contributors may be used to endorse or promote products
+ *         derived from this software without specific prior written permission.
+ *
+ *   THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+ *   AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+ *   IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE 
+ *   ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR CONTRIBUTORS BE
+ *   LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
+ *   CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
+ *   SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
+ *   INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
+ *   CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
+ *   ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
+ *   THE POSSIBILITY OF SUCH DAMAGE.
+ *****************************************************************************/
+#define PIX_SORT(a,b) { if ((a)>(b)) PIX_SWAP((a),(b)); }
+#define PIX_SWAP(a,b) { unsigned char temp=(a);(a)=(b);(b)=temp; }
+
+void IMG_median_3x3_8 (const unsigned char *restrict in_data, int cols, unsigned char * restrict out_data);
+void IMG_sobel_3x3_8  (const unsigned char *restrict in_data, unsigned char * restrict out_data, int rows, int cols);
+
+kernel void Median3x3(global const uchar* src, global uchar *dest,
+                      const int width, const int height,
+                      const int dstride, const int sstride)
+{
+  int i;
+  const int max_X = width - 1;
+  const int max_Y = height - 1;
+    for (i = 0; i < max_Y; i++) {
+      IMG_median_3x3_8 ((const unsigned char *)src, max_X, (const unsigned char *)dest);
+      src  += sstride;
+      dest += dstride; 
+    }
+}
+
+kernel void Sobel3x3(global const uchar* src, global uchar *dest,
+                      const int width, const int height,
+                      const int dstride, const int sstride)
+{
+  IMG_sobel_3x3_8((const unsigned char *)src, (const unsigned char *)dest, width, height);
+}
+
+kernel void Median2x2 (global const uchar* src, global uchar *dest,
+                      const int width, const int height,
+                      const int dstride, const int sstride)
+{
+/***
+    int id = get_global_id(0);
+    c[id] = a[id] + b[id];
+***/
+  unsigned char p[5];
+  int i, j, k;
+
+    /* copy the top and bottom rows into the result array */
+    for (i = 0; i < width; i++) {
+      dest[i] = src[i];
+      dest[(height - 1) * dstride + i] = src[(height - 1) * sstride + i];
+    }
+
+  /* process the interior pixels */
+  for (k = 2; k < height; k++) {
+    dest += dstride;
+    src += sstride;
+
+    dest[0] = src[0];
+    for (j = 2, i = 1; j < width; j++, i++) {
+      p[0] = src[i - sstride];
+      p[1] = src[i - 1];
+      p[2] = src[i];
+      p[3] = src[i + 1];
+      p[4] = src[i + sstride];
+      PIX_SORT (p[0], p[1]);
+      PIX_SORT (p[3], p[4]);
+      PIX_SORT (p[0], p[3]);
+      PIX_SORT (p[1], p[4]);
+      PIX_SORT (p[1], p[2]);
+      PIX_SORT (p[2], p[3]);
+      PIX_SORT (p[1], p[2]);
+      /* debug line */
+      if(k > 24 && k < 32) dest[i] = 0;
+      else
+      dest[i] = p[2];
+    }
+    dest[i] = src[i];
+  }
+}
+/* nothing past this point */

diff --git a/src/kernels/oclconv/oclconv.cpp b/src/kernels/oclconv/oclconv.cpp
new file mode 100644 (file)
index 0000000..7d070e1
--- /dev/null
+++ b/src/kernels/oclconv/oclconv.cpp
@@ -0,0 +1,152 @@
+/******************************************************************************
+ * Copyright (c) 2013-2014, Texas Instruments Incorporated - http://www.ti.com/
+ *   All rights reserved.
+ *
+ *   Redistribution and use in source and binary forms, with or without
+ *   modification, are permitted provided that the following conditions are met:
+ *       * Redistributions of source code must retain the above copyright
+ *         notice, this list of conditions and the following disclaimer.
+ *       * Redistributions in binary form must reproduce the above copyright
+ *         notice, this list of conditions and the following disclaimer in the
+ *         documentation and/or other materials provided with the distribution.
+ *       * Neither the name of Texas Instruments Incorporated nor the
+ *         names of its contributors may be used to endorse or promote products
+ *         derived from this software without specific prior written permission.
+ *
+ *   THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+ *   AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+ *   IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+ *   ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR CONTRIBUTORS BE
+ *   LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
+ *   CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
+ *   SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
+ *   INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
+ *   CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
+ *   ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
+ *   THE POSSIBILITY OF SUCH DAMAGE.
+ *****************************************************************************/
+#define __CL_ENABLE_EXCEPTIONS
+#include <CL/cl.hpp>
+#include <iostream>
+#include <fstream>
+#include <cstdlib>
+#include "ocl_util.h"
+#include "conv.dsp_h"
+
+//#define VERBOSE
+using namespace cl;
+using namespace std;
+
+static const int NumElements = 1920 * 1200; /* Maximum resolution */
+
+static cl_uchar src  [NumElements];
+static cl_uchar dst  [NumElements];
+
+static int oclconv_imgproc(char *kernelName, unsigned char *data_in, unsigned char *data_out, int width, int height, int sstride, int dstride)
+{
+   cl_int err     = CL_SUCCESS;
+   int  bufsize = sizeof(src);
+#ifdef VERBOSE
+   ofstream logfile;
+   logfile.open ("/home/root/oclconv_log.txt", ios::out | ios::app); 
+   logfile << "Entered oclconv_test, width=" << width << " height=" << height << " dstride=" << dstride << " sstride=" << sstride << '\n';
+   logfile.close();
+#endif
+   unsigned char *ptr_src     = src;
+   unsigned char *ptr_data_in = data_in;
+   for (int y=0; y < height; y++)
+   {
+     memcpy (ptr_src, ptr_data_in, width);
+     ptr_src     += sstride;
+     ptr_data_in += sstride;
+   }
+#ifdef VERBOSE
+   logfile.open ("/home/root/oclconv_log.txt", ios::out | ios::app); 
+   logfile << "Transfer done, starting openCL specific commands" << '\n';
+   logfile.close();
+#endif
+   try 
+   {
+     Context context(CL_DEVICE_TYPE_ACCELERATOR);
+     std::vector<Device> devices = context.getInfo<CL_CONTEXT_DEVICES>();
+     devices.resize(1); // resize to 1 since we are only running on 1 DSP
+
+     Buffer bufA   (context, CL_MEM_READ_ONLY,  bufsize);
+     Buffer bufDst (context, CL_MEM_WRITE_ONLY, bufsize);
+
+     Program::Binaries binary(1, make_pair(conv_dsp_bin,sizeof(conv_dsp_bin)));
+     Program           program = Program(context, devices, binary);
+     program.build(devices);
+     Kernel kernel(program, kernelName);
+     kernel.setArg(0, bufA);
+     kernel.setArg(1, bufDst);
+     kernel.setArg(2, width);
+     kernel.setArg(3, height);
+     kernel.setArg(4, dstride);
+     kernel.setArg(5, sstride);
+
+     Event ev1,ev2,ev3,ev4;
+
+     CommandQueue Q(context, devices[0], CL_QUEUE_PROFILING_ENABLE);
+
+     Q.enqueueWriteBuffer(bufA, CL_FALSE, 0, bufsize, src, NULL, &ev1);
+     //Q.enqueueNDRangeKernel(kernel, NullRange, NDRange(NumVecElements), 
+     //                       NDRange(WorkGroupSize), NULL, &ev3);
+     Q.enqueueTask (kernel, NULL, &ev3);
+     Q.enqueueReadBuffer (bufDst, CL_TRUE, 0, bufsize, dst, NULL, &ev4);
+   }
+   catch (Error err) 
+   { cerr << "ERROR: " << err.what() << "(" << err.err() << ")" << endl; }
+#ifdef VERBOSE
+   logfile.open ("/home/root/oclconv_log.txt", ios::out | ios::app); 
+   logfile << "OpenCL done, start transfer to output array!" << '\n';
+   logfile.close();
+#endif
+
+   unsigned char *ptr_dst      = dst;
+   unsigned char *ptr_data_out = data_out;
+   for (int y=0; y < height; y++)
+   {
+     memcpy (ptr_data_out, ptr_dst, width);
+     ptr_dst      += dstride;
+     ptr_data_out += dstride;
+   }
+
+#ifdef VERBOSE
+   logfile.open ("/home/root/oclconv_log.txt", ios::out | ios::app); 
+   logfile << "Success!" << endl; 
+   logfile.close();
+#endif
+   return 0;
+}
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+int oclconv_kernel(int kernel_type, int filter_size, 
+                   unsigned char *data_in, unsigned char *data_out, 
+                   int width, int height, int dstride, int sstride)
+{
+  int retval = -1;
+  switch(kernel_type)
+  {
+    case 0: /* Median */
+      if(filter_size == 5) { 
+        retval = oclconv_imgproc("Median2x2", data_in, data_out, width, height, sstride, dstride);
+      } else if(filter_size == 9) { 
+        retval = oclconv_imgproc("Median3x3", data_in, data_out, width, height, sstride, dstride);
+      }
+      break;
+    case 1: /* Sobel */
+      if(filter_size == 9) { 
+        retval = oclconv_imgproc("Sobel3x3", data_in, data_out, width, height, sstride, dstride);
+      }
+      break;
+    default:
+      break;
+  }
+  return retval;
+}
+#ifdef __cplusplus
+}
+#endif