mirror of https://github.com/opencv/opencv.git
salmanulhaq
14 years ago
parent
426c358352
commit
1e691daab4
6 changed files with 785 additions and 0 deletions
@ -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 <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);*/ |
||||||
|
|
||||||
|
} |
||||||
|
} |
||||||
|
} |
||||||
@ -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; |
||||||
|
} |
||||||
@ -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; } |
||||||
|
} |
||||||
|
} |
||||||
|
|
||||||
|
|
||||||
@ -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); |
||||||
|
|
||||||
|
} |
||||||
|
} |
||||||
|
} |
||||||
@ -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. */ |
||||||
@ -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 <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 |
||||||
Loading…
Reference in new issue