From 462d4a1dae1831f154cb2841ec00b699cd401018 Mon Sep 17 00:00:00 2001 From: Vadim Pisarevsky Date: Thu, 7 Jul 2011 16:40:44 +0000 Subject: [PATCH] opencl module is not ready for trunk yet --- modules/ocl/include/ocl.hpp | 141 ------------------ modules/ocl/include/ocl_util.h | 61 -------- modules/ocl/include/ocldefs.h | 59 -------- 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 -------- 9 files changed, 1046 deletions(-) delete mode 100644 modules/ocl/include/ocl.hpp delete mode 100644 modules/ocl/include/ocl_util.h delete mode 100644 modules/ocl/include/ocldefs.h delete mode 100644 modules/ocl/src/ocl/ocl_arithmetic.cpp delete mode 100644 modules/ocl/src/ocl/ocl_util.cpp delete mode 100644 modules/ocl/src/ocl/oclmat.cpp delete mode 100644 modules/ocl/src/ocl/opticalFlowLK.cpp delete mode 100644 modules/ocl/src/ocl/precomp.cpp delete mode 100644 modules/ocl/src/ocl/precomp.hpp diff --git a/modules/ocl/include/ocl.hpp b/modules/ocl/include/ocl.hpp deleted file mode 100644 index 82849b3..0000000 --- a/modules/ocl/include/ocl.hpp +++ /dev/null @@ -1,141 +0,0 @@ -/*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*/ - -#ifndef OCL_HPP -#define OCL_HPP - -#include "precomp.hpp" - - -//Root level namespace -namespace cv{ - - //This namespace will contain all the source code for the OpenCL module - namespace ocl{ - - extern cl_context ocl_context; - extern cl_command_queue ocl_cmd_queue; - - class OCL_EXPORTS OclMat{ - - public: - //! default constructor - OclMat(); - - //! constructs OclMat of the specified size and type (supported types are CV_8UC1, CV_8UC3, CV_16UC1, CV_16UC3, CV_64FC3 etc) - OclMat(int rows, int cols, int type); - OclMat(Size size, int type); - //! constucts OclMat and fills it with the specified value _s. - OclMat(int rows, int cols, int type, const Scalar& s); - OclMat(Size size, int type, const Scalar& s); - //! copy constructor - OclMat(const OclMat& m); - - //! builds OclMat from Mat. Perfom blocking upload to device. - explicit OclMat (const Mat& m); - - //! destructor - calls release() - ~OclMat(); - - //Releases the OpenCL context, command queue and the data buffer - void release(); - - //! assignment operators - OclMat& operator = (const OclMat& m); - //! assignment operator. Perfom blocking upload to device. - OclMat& operator = (const Mat& m); - - //! sets every OclMatelement to s - OclMat& operator = (const Scalar& s); - - //! sets some of the OclMat elements to s, according to the mask - OclMat& setTo(const Scalar& s); - - //! pefroms blocking upload data to GpuMat. - void upload(const cv::Mat& m); - - //! downloads data from device to host memory. Blocking calls. - operator Mat(); - void download(cv::Mat& m); - - //! returns the size of element in bytes. - size_t elemSize() const; - //! returns the size of element channel in bytes. - size_t elemSize1() const; - //! returns element type, similar to CV_MAT_TYPE(cvMat->type) - int type() const; - //! returns element type, similar to CV_MAT_DEPTH(cvMat->type) - int depth() const; - //! returns element type, similar to CV_MAT_CN(cvMat->type) - int channels() const; - //! returns step/elemSize1() - size_t step1() const; - //! returns GpuMatrix size: - // width == number of columns, height == number of rows - Size size() const; - //! returns true if OclMat data is NULL - bool empty() const; - - int flags; - - //! the number of rows and columns - int rows, cols; - //! a distance between successive rows in bytes; includes the gap if any - size_t step; - //! pointer to the data of type cl_mem - cl_mem data; - - int* refcount; - - //! allocates new OclMat data unless the OclMat already has specified size and type. - // previous data is unreferenced if needed. - void create(int rows, int cols, int type); - void create(Size size, int type); - - void _upload(size_t size, void* src); - void _download(size_t size, void* dst); - - }; - - //Creates the OpenCL context and command queue - OCL_EXPORTS void init(); - - } -} - -#endif \ No newline at end of file diff --git a/modules/ocl/include/ocl_util.h b/modules/ocl/include/ocl_util.h deleted file mode 100644 index bc85033..0000000 --- a/modules/ocl/include/ocl_util.h +++ /dev/null @@ -1,61 +0,0 @@ -/*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*/ - -#ifndef OCL_UTIL_H -#define OCL_UTIL_H - -#include "ocl.hpp" - -namespace cv{ - namespace ocl{ - namespace util{ - - //Step 1: Get the target platform - cl_platform_id GetOCLPlatform(); - - //Step 2: Create a context - int createContext(cl_context* context, cl_command_queue* cmd_queue, bool hasGPU); - } - } -} - -#endif \ No newline at end of file diff --git a/modules/ocl/include/ocldefs.h b/modules/ocl/include/ocldefs.h deleted file mode 100644 index 8afa4eb..0000000 --- a/modules/ocl/include/ocldefs.h +++ /dev/null @@ -1,59 +0,0 @@ -/*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*/ - -/* -This header file will contain all the type, struct and enum definitions exclusively used by the OpenCL module -*/ - -#ifndef OCLDEFS_H -#define OCLDEFS_H - -enum OclStatus{ - - OCL_SUCCESS=0, - OCL_ERROR=-100, - OCL_SIZE_ERROR = -101 - -}; - - -#endif \ No newline at end of file diff --git a/modules/ocl/src/ocl/ocl_arithmetic.cpp b/modules/ocl/src/ocl/ocl_arithmetic.cpp deleted file mode 100644 index a27e622..0000000 --- a/modules/ocl/src/ocl/ocl_arithmetic.cpp +++ /dev/null @@ -1,172 +0,0 @@ -/*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 deleted file mode 100644 index b592637..0000000 --- a/modules/ocl/src/ocl/ocl_util.cpp +++ /dev/null @@ -1,86 +0,0 @@ -/*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 deleted file mode 100644 index a27f92e..0000000 --- a/modules/ocl/src/ocl/oclmat.cpp +++ /dev/null @@ -1,253 +0,0 @@ -/*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 deleted file mode 100644 index f65eb36..0000000 --- a/modules/ocl/src/ocl/opticalFlowLK.cpp +++ /dev/null @@ -1,168 +0,0 @@ -/*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 deleted file mode 100644 index 2bf93e6..0000000 --- a/modules/ocl/src/ocl/precomp.cpp +++ /dev/null @@ -1,44 +0,0 @@ -/*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 deleted file mode 100644 index 6b1a562..0000000 --- a/modules/ocl/src/ocl/precomp.hpp +++ /dev/null @@ -1,62 +0,0 @@ -/*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