Chiebot-Mirror
/
opencv
mirror of https://github.com/opencv/opencv.git
8
0
Fork 0
Code Issues Projects Releases Wiki Activity

opencl module is not ready for trunk yet

Browse Source
pull/13383/head
Vadim Pisarevsky 14 years ago
parent 1e691daab4
commit 462d4a1dae
9 changed files with 0 additions and 1046 deletions
Whitespace
Show all changes Ignore whitespace when comparing lines Ignore changes in amount of whitespace Ignore changes in whitespace at EOL
Split View
Diff Options
Show Stats Download Patch File Download Diff File
  1. 141
      modules/ocl/include/ocl.hpp
  2. 61
      modules/ocl/include/ocl_util.h
  3. 59
      modules/ocl/include/ocldefs.h
  4. 172
      modules/ocl/src/ocl/ocl_arithmetic.cpp
  5. 86
      modules/ocl/src/ocl/ocl_util.cpp
  6. 253
      modules/ocl/src/ocl/oclmat.cpp
  7. 168
      modules/ocl/src/ocl/opticalFlowLK.cpp
  8. 44
      modules/ocl/src/ocl/precomp.cpp
  9. 62
      modules/ocl/src/ocl/precomp.hpp

141
modules/ocl/include/ocl.hpp
Unescape View File

@ -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

61
modules/ocl/include/ocl_util.h
Unescape View File

@ -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

59
modules/ocl/include/ocldefs.h
Unescape View File

@ -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

172
modules/ocl/src/ocl/ocl_arithmetic.cpp
Unescape View File

@ -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 <iostream>
#include <fstream>
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);*/
}
}
}

86
modules/ocl/src/ocl/ocl_util.cpp
Unescape View File

@ -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;
}

253
modules/ocl/src/ocl/oclmat.cpp
Unescape View File

@ -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; }
}
}

168
modules/ocl/src/ocl/opticalFlowLK.cpp
Unescape View File

@ -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);
}
}
}

44
modules/ocl/src/ocl/precomp.cpp
Unescape View File

@ -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. */

62
modules/ocl/src/ocl/precomp.hpp
Unescape View File

@ -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 <exception>
#include "opencv/highgui.h"
#include "opencv/cv.h"
#include <CL/cl.h> //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 <stdio.h>
#endif
#endif
Write Preview
Loading…
Cancel
Save