summary | shortlog | log | commit | commitdiff | tree
raw | patch | inline | side by side (parent: 7808287)
raw | patch | inline | side by side (parent: 7808287)
author | Djordje Senicic <d-senicic1@ti.com> | |
Sat, 27 Feb 2016 01:21:43 +0000 (20:21 -0500) | ||
committer | Djordje Senicic <d-senicic1@ti.com> | |
Sat, 27 Feb 2016 01:21:43 +0000 (20:21 -0500) |
src/kernels/make.inc | [new file with mode: 0644] | patch | blob |
src/kernels/oclconv/conv.cl | [new file with mode: 0644] | patch | blob |
src/kernels/oclconv/oclconv.cpp | [new file with mode: 0644] | patch | blob |
diff --git a/src/kernels/make.inc b/src/kernels/make.inc
--- /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
--- /dev/null
@@ -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
--- /dev/null
@@ -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