From 1e691daab430b5fa4082c952e84ad49cdf3c53d3 Mon Sep 17 00:00:00 2001 From: salmanulhaq Date: Thu, 7 Jul 2011 16:35:20 +0000 Subject: [PATCH] ocl folder contains the gpu implementations --- modules/ocl/src/ocl/ocl_arithmetic.cpp | 172 ++++++++++++++++++++++ modules/ocl/src/ocl/ocl_util.cpp | 86 +++++++++++ modules/ocl/src/ocl/oclmat.cpp | 253 +++++++++++++++++++++++++++++++++ modules/ocl/src/ocl/opticalFlowLK.cpp | 168 ++++++++++++++++++++++ modules/ocl/src/ocl/precomp.cpp | 44 ++++++ modules/ocl/src/ocl/precomp.hpp | 62 ++++++++ 6 files changed, 785 insertions(+) create mode 100644 modules/ocl/src/ocl/ocl_arithmetic.cpp create mode 100644 modules/ocl/src/ocl/ocl_util.cpp create mode 100644 modules/ocl/src/ocl/oclmat.cpp create mode 100644 modules/ocl/src/ocl/opticalFlowLK.cpp create mode 100644 modules/ocl/src/ocl/precomp.cpp create mode 100644 modules/ocl/src/ocl/precomp.hpp diff --git a/modules/ocl/src/ocl/ocl_arithmetic.cpp b/modules/ocl/src/ocl/ocl_arithmetic.cpp new file mode 100644 index 0000000..a27e622 --- /dev/null +++ b/modules/ocl/src/ocl/ocl_arithmetic.cpp @@ -0,0 +1,172 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2000-2008, Intel Corporation, all rights reserved. +// Copyright (C) 2009, Willow Garage Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * The name of the copyright holders may not 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 Intel Corporation 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. +// +//M*/ + + +#include "ocl.hpp" +#include +#include +using namespace std; + +namespace cv{ + namespace ocl{ + + void writeBinaries(cl_program cpProgram) + { + ofstream myfile("kernel.ptx"); + + cl_uint program_num_devices = 1; + clGetProgramInfo(cpProgram, CL_PROGRAM_NUM_DEVICES, sizeof(cl_uint), &program_num_devices, NULL); + + if (program_num_devices == 0) + return; + + size_t binaries_sizes[1]; + + clGetProgramInfo(cpProgram, CL_PROGRAM_BINARY_SIZES, program_num_devices*sizeof(size_t), binaries_sizes, NULL); + + char **binaries = new char*[1]; + binaries[0] = 0; + + for (size_t i = 0; i < 1; i++) + binaries[i] = new char[binaries_sizes[i]+1]; + + + clGetProgramInfo(cpProgram, CL_PROGRAM_BINARIES, program_num_devices*sizeof(size_t), binaries, NULL); + + if(myfile.is_open()) + { + for (size_t i = 0; i < program_num_devices; i++) + { + myfile << binaries[i]; + } + } + myfile.close(); + + for (size_t i = 0; i < program_num_devices; i++) + delete [] binaries; + + delete [] binaries; +} + + OCL_EXPORTS void add(const OclMat& a, const OclMat& b, OclMat& sum){ + + if(a.rows != b.rows || a.cols != b.cols) + return; + + int size = a.rows*a.cols; + + cl_program program; + cl_kernel kernel; + size_t global_work_size[1]; + size_t local_work_size[1]; + cl_uint size_ret = 0; + cl_int err; + + const char* add_kernel_source[] = {\ + "__kernel void add_kernel (__global const uchar *a, __global const uchar* b, __global uchar* c)"\ + "{"\ + "int tid = get_global_id(0);"\ + "c[tid] = a[tid] + b[tid];"\ + "}" + }; + + program = clCreateProgramWithSource(ocl_context, 1, (const char**)&add_kernel_source, NULL, NULL); +/////////////////// + writeBinaries(program); +/////////////////// + + + //err = clBuildProgram(program, 0, NULL, NULL, NULL, NULL); + /* FILE* fp = fopen("kernel.ptx", "r"); + fseek (fp , 0 , SEEK_END); + const size_t lSize = ftell(fp); + rewind(fp); + unsigned char* buffer; + buffer = (unsigned char*) malloc (lSize); + fread(buffer, 1, lSize, fp); + fclose(fp); + + size_t cb; + err = clGetContextInfo(ocl_context, CL_CONTEXT_DEVICES, 0, NULL, &cb); + cl_device_id *devices = (cl_device_id*)malloc(cb); + clGetContextInfo(ocl_context, CL_CONTEXT_DEVICES, cb, devices, NULL); + + cl_int status; + program = clCreateProgramWithBinary(ocl_context, 1, devices, + &lSize, (const unsigned char**)&buffer, + &status, &err); + + err = clBuildProgram(program, 0, NULL, NULL, NULL, NULL); +#ifdef _DEBUG + if(err != CL_SUCCESS){ + printf("(Error code: %d)Build failed, check the program source...\n",err); + return; + } +#endif + + //writeBinaries(program); + + kernel = clCreateKernel(program, "add_kernel", NULL); + + err = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &a.data); + err = clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &a.data); + err = clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &a.data); +#ifdef _DEBUG + if(err != CL_SUCCESS){ + printf("(Error code: %d)Failed at setting kernel arguments...\n",err); + return; + } +#endif + + global_work_size[0] = size; + local_work_size[0]= 1; + + err = clEnqueueNDRangeKernel(ocl_cmd_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL); +#ifdef _DEBUG + if(err != CL_SUCCESS){ + printf("(Error code: %d)Kernel execution failed...\n",err); + return; + } +#endif + + clReleaseKernel(kernel); + clReleaseProgram(program);*/ + + } + } +} \ No newline at end of file diff --git a/modules/ocl/src/ocl/ocl_util.cpp b/modules/ocl/src/ocl/ocl_util.cpp new file mode 100644 index 0000000..b592637 --- /dev/null +++ b/modules/ocl/src/ocl/ocl_util.cpp @@ -0,0 +1,86 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2000-2008, Intel Corporation, all rights reserved. +// Copyright (C) 2009, Willow Garage Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * The name of the copyright holders may not 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 Intel Corporation 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. +// +//M*/ + +#include "ocl.hpp" + +cl_platform_id cv::ocl::util::GetOCLPlatform() +{ + cl_platform_id pPlatforms[10] = { 0 }; + char pPlatformName[128] = { 0 }; + + cl_uint uiPlatformsCount = 0; + cl_int err = clGetPlatformIDs(10, pPlatforms, &uiPlatformsCount); + for (cl_uint ui = 0; ui < uiPlatformsCount; ++ui) + { + err = clGetPlatformInfo(pPlatforms[ui], CL_PLATFORM_NAME, 128 * sizeof(char), pPlatformName, NULL); + if ( err != CL_SUCCESS ) + { + return NULL; + } + + return pPlatforms[ui]; + } +} + +int cv::ocl::util::createContext(cl_context* context, cl_command_queue* cmd_queue, bool hasGPU){ + + size_t cb; + cl_int err; + cl_platform_id platform_id = cv::ocl::util::GetOCLPlatform(); + + cl_context_properties context_properties[3] = {CL_CONTEXT_PLATFORM, (cl_context_properties)platform_id, NULL }; + + //If the GPU is present, create the context on GPU + if(hasGPU) + *context = clCreateContextFromType(context_properties, CL_DEVICE_TYPE_GPU, NULL, NULL, NULL); + else + *context = clCreateContextFromType(context_properties, CL_DEVICE_TYPE_CPU, NULL, NULL, NULL); + + // get the list of devices associated with context + err = clGetContextInfo(*context, CL_CONTEXT_DEVICES, 0, NULL, &cb); + + cl_device_id *devices = (cl_device_id*)malloc(cb); + + clGetContextInfo(*context, CL_CONTEXT_DEVICES, cb, devices, NULL); + + // create a command-queue + *cmd_queue = clCreateCommandQueue(*context, devices[0], 0, NULL); + free(devices); + + return CL_SUCCESS; +} diff --git a/modules/ocl/src/ocl/oclmat.cpp b/modules/ocl/src/ocl/oclmat.cpp new file mode 100644 index 0000000..a27f92e --- /dev/null +++ b/modules/ocl/src/ocl/oclmat.cpp @@ -0,0 +1,253 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2000-2008, Intel Corporation, all rights reserved. +// Copyright (C) 2009, Willow Garage Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other GpuMaterials provided with the distribution. +// +// * The name of the copyright holders may not 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 Intel Corporation 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. +// +//M*/ + +#include "ocl.hpp" + +namespace cv{ + + namespace ocl{ + + cl_context ocl_context; + cl_command_queue ocl_cmd_queue; + bool initialized = false; + + inline OclMat::OclMat() : rows(0), cols(0), step(0), data(0), refcount(0) {} + + inline OclMat::OclMat(int _rows, int _cols, int _type): flags(0), rows(0), cols(0), step(0), data(0), refcount(0){ + + if(_rows > 0 && _cols > 0) + create(_rows, _cols, _type); + } + + inline OclMat::OclMat(Size size, int _type): flags(0), rows(0), cols(0), step(0), data(0), refcount(0){ + + if(size.height > 0 && size.width > 0) + create(size, _type); + } + + inline OclMat::OclMat(const OclMat& m) + : flags(m.flags), rows(m.rows), cols(m.cols), step(m.step), data(m.data), refcount(m.refcount){ + + if(refcount) + CV_XADD(refcount, 1); + } + + inline OclMat::OclMat(const Mat& m) + : flags(0), rows(0), cols(0), step(0), data(0), refcount(0) { upload(m); } + + inline OclMat::~OclMat(){ release(); } + + void OclMat::_upload(size_t size, void* src){ + + this->data = clCreateBuffer(ocl_context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, size, src, NULL); + } + + void OclMat::_download(size_t size, void* dst){ + + cl_int err = clEnqueueReadBuffer(ocl_cmd_queue, data, CL_TRUE, 0, size, dst, 0, NULL, NULL); + } + + + void OclMat::release(){ + + if( refcount && CV_XADD(refcount, -1) == 1 ) + { + free(refcount); + clReleaseMemObject(data); + } + clReleaseMemObject(data); + data = 0; + step = rows = cols = 0; + refcount = 0; + } + + void OclMat::upload(const Mat& m){ + + create(m.rows, m.cols, m.type()); + int ch = channels(); + int d = elemSize(); + size_t s = rows*cols*ch*d; + if(initialized) + this->_upload(s, m.data); + else{ + init(); + this->_upload(s, m.data); +#ifdef _DEBUG + printf("Context and Command queues not initialized. First call cv::ocl::init()"); +#endif + } + } + + void OclMat::download(Mat& m){ + + size_t s = rows*cols*channels()*elemSize(); + m.create(rows, cols, type()); + if(initialized){ + this->_download(s, m.data); + } + + else{ + init(); + this->_download(s, m.data); +#ifdef _DEBUG + printf("Context and Command queues not initialized. First call cv::ocl::init()"); +#endif + + } + + } + + void init(){ + + if(!initialized){ + cv::ocl::util::createContext(&ocl_context, &ocl_cmd_queue, false); + initialized = true; + } + } + + void OclMat::create(int _rows, int _cols, int _type){ + + if(!initialized) + init(); + + _type &= TYPE_MASK; + + if( rows == _rows && cols == _cols && type() == _type && data ) + return; + if( data ) + release(); + + if( _rows > 0 && _cols > 0 ){ + flags = Mat::MAGIC_VAL + _type; + rows = _rows; + cols = _cols; + + size_t esz = elemSize(); + int ch = channels(); + + step = esz*cols*ch; + + size_t size = esz*rows*cols*ch; + data = clCreateBuffer(ocl_context, CL_MEM_READ_WRITE, size, NULL, NULL); + + if (esz * cols == step) + flags |= Mat::CONTINUOUS_FLAG; + + /* + refcount = (int*)fastMalloc(sizeof(*refcount)); + *refcount = 1; + */ + } + } + + void OclMat::create(Size size, int _type){ + + return create(size.height, size.width, _type); + } + + inline OclMat::operator Mat() + { + Mat m; + download(m); + return m; + } + + inline OclMat& OclMat::operator = (const OclMat& m) + { + if( this != &m ) + { + if( m.refcount ) + CV_XADD(m.refcount, 1); + release(); + flags = m.flags; + rows = m.rows; cols = m.cols; + step = m.step; data = m.data; + data = m.data; + refcount = m.refcount; + } + return *this; + } + + inline OclMat::OclMat(int _rows, int _cols, int _type, const Scalar& _s) + : flags(0), rows(0), cols(0), step(0), data(0), refcount(0) + { + if(_rows > 0 && _cols > 0) + { + create(_rows, _cols, _type); + *this = _s; + } + } + + inline OclMat& OclMat::operator = (const Mat& m) { upload(m); return *this; } + + OclMat& OclMat::operator = (const Scalar& s) + { + setTo(s); + return *this; + } + + OclMat& OclMat::setTo(const Scalar& s){ + + //if (s[0] == 0.0 && s[1] == 0.0 && s[2] == 0.0 && s[3] == 0.0) + //{ + + size_t sz = rows*cols*channels()*elemSize(); + //void* ptr = (void*)malloc(sz); + //memset(ptr, s[0], sz); + //clEnqueueWriteBuffer(ocl_cmd_queue, data, CL_TRUE, 0, sz, ptr, NULL, NULL, NULL); + //free(ptr); ptr = 0; + return *this; + //} + } + + inline size_t OclMat::elemSize() const{ return CV_ELEM_SIZE(flags); } + inline size_t OclMat::elemSize1() const{ return CV_ELEM_SIZE1(flags); } + inline int OclMat::type() const { return CV_MAT_TYPE(flags); } + inline int OclMat::depth() const{ return CV_MAT_DEPTH(flags); } + inline int OclMat::channels() const{ return CV_MAT_CN(flags); } + inline size_t OclMat::step1() const{ return step/elemSize1(); } + inline Size OclMat::size() const{ return Size(cols, rows); } + inline bool OclMat::empty() const{ return data == 0; } + } +} + + \ No newline at end of file diff --git a/modules/ocl/src/ocl/opticalFlowLK.cpp b/modules/ocl/src/ocl/opticalFlowLK.cpp new file mode 100644 index 0000000..f65eb36 --- /dev/null +++ b/modules/ocl/src/ocl/opticalFlowLK.cpp @@ -0,0 +1,168 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2000-2008, Intel Corporation, all rights reserved. +// Copyright (C) 2009, Willow Garage Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other GpuMaterials provided with the distribution. +// +// * The name of the copyright holders may not 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 Intel Corporation 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. +// +//M*/ + +#include "ocl.hpp" + +namespace cv{ + namespace ocl{ + + extern cl_context ocl_context; + extern cl_command_queue ocl_cmd_queue; + + OCL_EXPORTS void opticalFlowLK(const OclMat& imgA, const OclMat& imgB, CvSize winSize, OclMat& velX, OclMat& velY){ + + if(imgA.rows != imgB.rows || imgA.cols != imgB.cols) + return; + + int size = imgA.rows*imgA.cols; + + cl_program program; + cl_kernel kernel; + size_t global_work_size[1]; + size_t local_work_size[1]; + cl_uint size_ret = 0; + cl_int err; + + + //, __global const ushort* imWidth, __global const ushort* hradius + const char* of_kernel_source[] = {\ + "__kernel void derivatives (__global const uchar *imgA, __global const uchar* imgB, __global float* fx, __global float* fy, __global float* ft, __global float* u, __global float* v)"\ + "{"\ + "int tid = get_global_id(0);"\ + "float Ix = 0.0f, Iy = 0.0f, It = 0.0f;"\ + "ushort imageWidth = 1024;"\ + "ushort half_radius = 1;"\ + "Ix = ((imgA[tid+1] - imgA[tid] + imgB[tid+1] - imgB[tid] )/2);"\ + "Iy = ((imgA[tid + imageWidth] - imgA[tid] + imgB[tid + imageWidth] - imgB[tid])/2);" + "It = imgB[tid] - imgA[tid];"\ + "fx[tid] = Ix;"\ + "fy[tid] = Iy;"\ + "ft[tid] = It;"\ + "__local float s_data[3];"\ + "float A = 0.0f, B = 0.0f, C = 0.0f, D = 0.0f, E = 0.0f;"\ + "short i = 0;"\ + "for (i = -half_radius; i <= half_radius; i++){"\ + "for (short j = -half_radius; j <= half_radius; j++){"\ + "s_data[0] = fx[tid + i + j*imageWidth];"\ + "s_data[1] = fy[tid + i + j*imageWidth];"\ + "s_data[2] = ft[tid + i + j*imageWidth];"\ + "A = A + powf(s_data[0],2);"\ + "B = B + powf(s_data[1],2);"\ + "C = C + s_data[0] * s_data[1];"\ + "D = D + s_data[0] * s_data[2];"\ + "E = E + s_data[1] * s_data[2];"\ + "}}"\ + "u[tid] = (D*B - E*C)/(A*B - C*C);"\ + "v[tid] = (E*A - D*C)/(A*B - C*C);"\ + "}" + }; + + //program = clCreateProgramWithSource(imgA.ocl_context, 1, (const char**)&of_kernel_source, NULL, NULL); + program = clCreateProgramWithSource(ocl_context, 1, (const char**)&of_kernel_source, NULL, NULL); + + err = clBuildProgram(program, 0, NULL, NULL, NULL, NULL); + +#ifdef _DEBUG + if(err != CL_SUCCESS){ + printf("(Error code: %d)Build failed, check the program source...\n",err); + return; + } +#endif + + kernel = clCreateKernel(program, "derivatives", NULL); + + //Creating additional temporary buffers fx, fy and ft + //cl_mem fx = clCreateBuffer(imgA.ocl_context, CL_MEM_READ_WRITE, size*sizeof(cl_float), NULL, NULL); + //cl_mem fy = clCreateBuffer(imgA.ocl_context, CL_MEM_READ_WRITE, size*sizeof(cl_float), NULL, NULL); + //cl_mem ft = clCreateBuffer(imgA.ocl_context, CL_MEM_READ_WRITE, size*sizeof(cl_float), NULL, NULL); + + cl_mem fx = clCreateBuffer(ocl_context, CL_MEM_READ_WRITE, size*sizeof(cl_float), NULL, NULL); + cl_mem fy = clCreateBuffer(ocl_context, CL_MEM_READ_WRITE, size*sizeof(cl_float), NULL, NULL); + cl_mem ft = clCreateBuffer(ocl_context, CL_MEM_READ_WRITE, size*sizeof(cl_float), NULL, NULL); + + //Creating variables for imageWidth and half_window + ushort x_radius = winSize.width/2; + ushort cols = (ushort)imgA.cols; + + //cl_mem imageWidth = clCreateBuffer(ocl_context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeof(cl_ushort), &cols, NULL); + //cl_mem half_radius = clCreateBuffer(ocl_context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeof(cl_ushort), &x_radius, NULL); + + + err = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &imgA.data); + err = clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &imgB.data); + err = clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &fx); + err = clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &fy); + err = clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *) &ft); + err = clSetKernelArg(kernel, 5, sizeof(cl_mem), (void *) &velX.data); + err = clSetKernelArg(kernel, 6, sizeof(cl_mem), (void *) &velY.data); + //err = clSetKernelArg(kernel, 7, sizeof(cl_mem), (void *) &imageWidth); + //err = clSetKernelArg(kernel, 8, sizeof(cl_mem), (void *) &half_radius); + +#ifdef _DEBUG + if(err != CL_SUCCESS){ + printf("(Error code: %d)Failed at setting kernel arguments...\n",err); + return; + } +#endif + + global_work_size[0] = size; + local_work_size[0]= 1; + + //err = clEnqueueNDRangeKernel(imgA.ocl_cmd_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL); + err = clEnqueueNDRangeKernel(ocl_cmd_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL); +#ifdef _DEBUG + if(err != CL_SUCCESS){ + printf("(Error code: %d)Kernel execution failed...\n",err); + return; + } +#endif + clReleaseMemObject(fx); + clReleaseMemObject(fy); + clReleaseMemObject(ft); + //clReleaseMemObject(imageWidth); + //clReleaseMemObject(half_radius); + clReleaseKernel(kernel); + clReleaseProgram(program); + + } + } +} diff --git a/modules/ocl/src/ocl/precomp.cpp b/modules/ocl/src/ocl/precomp.cpp new file mode 100644 index 0000000..2bf93e6 --- /dev/null +++ b/modules/ocl/src/ocl/precomp.cpp @@ -0,0 +1,44 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// Intel License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2000, Intel Corporation, all rights reserved. +// Third party copyrights are property of their respective owners. +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's 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. +// +// * The name of Intel Corporation may not 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 Intel Corporation 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. +// +//M*/ + +#include "precomp.hpp" + +/* End of file. */ \ No newline at end of file diff --git a/modules/ocl/src/ocl/precomp.hpp b/modules/ocl/src/ocl/precomp.hpp new file mode 100644 index 0000000..6b1a562 --- /dev/null +++ b/modules/ocl/src/ocl/precomp.hpp @@ -0,0 +1,62 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2000-2008, Intel Corporation, all rights reserved. +// Copyright (C) 2009, Willow Garage Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's 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. +// +// * The name of the copyright holders may not 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 Intel Corporation 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. +// +//M*/ + +#ifndef PRECOMP_H__ +#define PRECOMP_H__ + +#define OCL_EXPORTS __declspec(dllexport) + +#include + +#include "opencv/highgui.h" +#include "opencv/cv.h" +#include //General OpenCL include file. Current version is written for Intel architecture using Intel's OpenCL SDK +#include "ocl.hpp" +#include "oclmat.hpp" +#include "ocldefs.h" +#include "ocl_util.h" + +#ifdef _DEBUG + #include +#endif + +#endif \ No newline at end of file -- 2.7.4