public/feature-extraction-opencl
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

init

Browse Source
This commit is contained in:
gfuse
2021-06-18 14:57:59 +02:00
commit 6e6d191bda
19 changed files with 1797 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files
CMakeLists.txt
Executable
+14
View File
@@ -0,0 +1,14 @@
cmake_minimum_required(VERSION 3.17)
project(Vio_OpenCL)
set(CMAKE_CXX_STANDARD 14)
find_package (OpenCL REQUIRED)
find_package(OpenCV REQUIRED)
include_directories (include ${OpenCL_INCLUDE_DIRS} ${OpenCV_INCLUDE_DIRS} kernel)
add_executable(${PROJECT_NAME} main.cpp include/cl_class.h src/cl_class.cpp)
target_compile_features(${PROJECT_NAME} PRIVATE cxx_range_for)
target_link_libraries(${PROJECT_NAME} ${OpenCL_LIBRARIES} ${OpenCV_LIBS})
ADD_DEFINITIONS(-DTEST_DATA_DIR=\"${PROJECT_SOURCE_DIR}/test/data\")
frame0287.jpg
Executable
BIN
View File
Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 59 KiB

include/FastGrayStep.hh
+80
View File
@@ -0,0 +1,80 @@
// Copyright (C) 2011 Dmitri Nikulin
// Copyright (C) 2011 Monash University
//
// Permission is hereby granted, free of charge, to any person
// obtaining a copy of this software and associated documentation
// files (the "Software"), to deal in the Software without
// restriction, including without limitation the rights to use,
// copy, modify, merge, publish, distribute, sublicense, and/or sell
// copies of the Software, and to permit persons to whom the
// Software is furnished to do so, subject to the following
// conditions:
//
// The above copyright notice and this permission notice shall be
// included in all copies or substantial portions of the Software.
//
// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
// EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES
// OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND
// NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT
// HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY,
// WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
// FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR
// OTHER DEALINGS IN THE SOFTWARE.
#ifndef __CVD_CL_FAST_GRAY_STEP_HH__
#define __CVD_CL_FAST_GRAY_STEP_HH__
#include <cvd-cl/states/ImageState.hh>
#include <cvd-cl/states/ListState.hh>
#include <cvd-cl/worker/WorkerStep.hh>
namespace CVD {
namespace CL {
/// \brief WorkerStep to filter a point list by FAST corner criteria.
class FastGrayStep : public WorkerStep {
public:
/// \brief Construct the step.
///
/// \param i_image Input image.
/// \param i_points Input point list.
/// \param o_points Output point list.
/// \param threshold FAST threshold.
/// \param ring Consecutive bits required in FAST ring.
explicit FastGrayStep(GrayImageState & i_image, PointListState & i_points, PointListState & o_points, cl_int threshold=40, cl_int ring=9);
/// \brief De-construct the step.
virtual ~FastGrayStep();
virtual void execute();
protected:
/// \brief Input image.
GrayImageState & i_image;
/// \brief Input point list.
PointListState & i_points;
/// \brief Output point list.
PointListState & o_points;
/// \brief FAST threshold.
cl_int const threshold;
/// \brief Consecutive bits required in FAST ring.
cl_int const ring;
/// \brief OpenCL program.
cl::Program program;
/// \brief OpenCL kernel.
cl::Kernel kernel;
};
} // namespace CL
} // namespace CVD
#endif /* __CVD_CL_FAST_GRAY_STEP_HH__ */
include/PreFastGrayStep.hh
+72
View File
@@ -0,0 +1,72 @@
// Copyright (C) 2011 Dmitri Nikulin
// Copyright (C) 2011 Monash University
//
// Permission is hereby granted, free of charge, to any person
// obtaining a copy of this software and associated documentation
// files (the "Software"), to deal in the Software without
// restriction, including without limitation the rights to use,
// copy, modify, merge, publish, distribute, sublicense, and/or sell
// copies of the Software, and to permit persons to whom the
// Software is furnished to do so, subject to the following
// conditions:
//
// The above copyright notice and this permission notice shall be
// included in all copies or substantial portions of the Software.
//
// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
// EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES
// OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND
// NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT
// HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY,
// WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
// FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR
// OTHER DEALINGS IN THE SOFTWARE.
#ifndef __CVD_CL_PRE_FAST_GRAY_STEP_HH__
#define __CVD_CL_PRE_FAST_GRAY_STEP_HH__
#include <cvd-cl/worker/WorkerStep.hh>
#include <cvd-cl/states/ImageState.hh>
#include <cvd-cl/states/ListState.hh>
namespace CVD {
namespace CL {
/// \brief WorkerStep to select points that are reasonably likely to later satisfy a FAST corner test.
class PreFastGrayStep : public WorkerStep {
public:
/// \brief Construct the step.
///
/// \param i_image Input image.
/// \param o_points Output point list.
/// \param threshold FAST threshold.
explicit PreFastGrayStep(GrayImageState & i_image, PointListState & o_points, cl_int threshold=40);
/// \brief De-construct the step.
virtual ~PreFastGrayStep();
virtual void execute();
protected:
/// \brief Input image.
GrayImageState & i_image;
/// \brief Output point list.
PointListState & o_points;
/// \brief FAST threshold.
cl_int const threshold;
/// \brief OpenCL program.
cl::Program program;
/// \brief OpenCL kernel.
cl::Kernel kernel;
};
} // namespace CL
} // namespace CVD
#endif /* __CVD_CL_PRE_FAST_GRAY_STEP_HH__ */
include/cl_class.h
+189
View File
@@ -0,0 +1,189 @@
//
// Created by root on 4/27/21.
//
#ifndef VIO_OPENCL_CL_CLASS_H
#define VIO_OPENCL_CL_CLASS_H
#include <CL/cl.hpp>
#include <iostream>
#include <vector>
#include <algorithm>
#include <fstream>
#include <exception>
#include <CL/opencl.h>
#include <opencv2/opencv.hpp>
class kernel{
public:
kernel(cl::Program* program,std::string name){_kernel = new cl::Kernel(*program,name.c_str());};
~kernel(){
_buffers.clear();
_images.clear();
};
template<typename T>
int32_t write(size_t id/*buffer ID*/,T* buf,cl::CommandQueue* queue,cl::Context* context,size_t buf_size){
_buffers.push_back(std::pair<std::pair<cl::Buffer,size_t>,size_t>(std::pair<cl::Buffer,size_t>(cl::Buffer(*context, CL_MEM_READ_WRITE, sizeof(T) * buf_size),
buf_size),id));
cl_int err=queue->enqueueWriteBuffer(_buffers.back().first.first,CL_TRUE, 0, sizeof(T) * buf_size, buf);
if(err!=0){
std::cerr<<"Error C:Kernel, F:write, L:enqueueWriteBuffer\n"<<err<<'\n';
return err;
}
err=_kernel->setArg(id,_buffers.back().first.first);
if(err!=0){
std::cerr<<"Error C:Kernel, F:write, L3:setArg\n"<<err<<'\n';
return err;
}
return CL_SUCCESS;
};
int32_t write(size_t id/*buffer ID*/,cv::Mat& buf,cl::Context* context){
_images.push_back(std::pair<cl::Image2D,size_t>(cl::Image2D(*context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
cl::ImageFormat(CL_R, CL_UNSIGNED_INT8),
buf.size().width,
buf.size().height,
0,
reinterpret_cast<uchar*>(buf.data)),id));
cl_int err=_kernel->setArg(id,_images.back().first);
if(err!=0){
std::cerr<<"Error C:Kernel, F:write image, L:setArg\n"<<err<<'\n';
return err;
}
return CL_SUCCESS;
};
template<typename T>
int32_t reload(size_t id,T* buf,cl::CommandQueue* queue){
try{
if(id>_buffers.size())return CL_MAP_FAILURE;
for(auto i:_buffers)if(i.second==id)
queue->enqueueWriteBuffer(i.first.first,CL_TRUE, 0, sizeof(T) * i.first.second, buf);
return CL_SUCCESS;
}catch (std::exception& err){
std::cerr<<"Reload buffer on GPU failed\n"<<err.what()<<'\n';
return CL_MAP_FAILURE;
}
};
int32_t reload(size_t id,cv::Mat& buf,cl::Context* context){
try{
for(auto i:_images)if(i.second==id)
i.first=cl::Image2D(*context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
cl::ImageFormat(CL_R, CL_UNSIGNED_INT8),
buf.size().width,
buf.size().height,
0,
reinterpret_cast<uchar*>(buf.data));
return CL_SUCCESS;
}catch (std::exception& err){
std::cerr<<"Reload buffer on GPU failed\n"<<err.what()<<'\n';
return CL_MAP_FAILURE;
}
};
cl::Kernel* _kernel = NULL;
std::pair<cl::Buffer,size_t> read(size_t id){
for(auto i:_buffers)if(i.second==id)return i.first;
}
cl::Image2D read_img(size_t id){
for(auto i:_images)if(i.second==id)return i.first;
}
private:
std::vector<std::pair<std::pair<cl::Buffer,size_t>,size_t>> _buffers;
std::vector<std::pair<cl::Image2D,size_t>> _images;
};
class read_cl{
public:
read_cl(std::string path){
src_str=(char *)malloc(0x100000);
FILE *fp;
fp = fopen(path.c_str(), "r");
size = fread(src_str, 1, 0x100000, fp);
fclose(fp);
}
void print(){
std::cerr<<src_str<<"\n";
}
char *src_str = nullptr;
size_t size;
};
class opencl{
public:
opencl();
~opencl();
int32_t make_kernel(std::string name){_kernels.push_back(kernel(program,name));};
template<typename T>
int32_t write_buf(size_t id1/*kernal ID*/,size_t id2/*buffer ID*/, int size,T* buf) {
try{
_kernels.at(id1).write(id2,buf,queue,context,size);
return CL_SUCCESS;
}catch (std::exception& err){
std::cerr<<"Write buffer on GPU failed\n"<<err.what()<<'\n';
return CL_MAP_FAILURE;
}
}
int32_t write_buf(size_t id1/*kernal ID*/,size_t id2/*buffer ID*/,cv::Mat& buf) {
try{
_kernels.at(id1).write(id2,buf,context);
return CL_SUCCESS;
}catch (std::exception& err){
std::cerr<<"Write buffer on GPU failed\n"<<err.what()<<'\n';
return CL_MAP_FAILURE;
}
}
template<typename T>
int32_t write_buf(size_t id1/*kernal ID*/,size_t id2/*buffer ID*/,T&& value){
_kernels.at(id1)._kernel->setArg(id2,value);
};
template<typename T>
int32_t reload_buf(size_t id1/*kernal ID*/,size_t id2/*buffer ID*/,T* buf) {
try{
_kernels.at(id1).reload(id2,buf,queue);
return CL_SUCCESS;
}catch (std::exception& err){
std::cerr<<"Reload buffer on GPU failed\n"<<err.what()<<'\n';
return CL_MAP_FAILURE;
}
}
int32_t reload_buf(size_t id1/*kernal ID*/,size_t id2/*buffer ID*/,cv::Mat& buf) {
try{
_kernels.at(id1).reload(id2,buf,context);
return CL_SUCCESS;
}catch (std::exception& err){
std::cerr<<"Write buffer on GPU failed\n"<<err.what()<<'\n';
return CL_MAP_FAILURE;
}
}
int32_t run(size_t id1/*kernal ID*/,std::size_t x=1,std::size_t y=1,std::size_t z=1) {
cl_int err=0;
if(z>1 && y>1){
err=queue->enqueueNDRangeKernel(*_kernels.at(id1)._kernel, cl::NullRange/*offset*/, cl::NDRange(x,y,z)/*Global*/, cl::NullRange/*local*/);
}else if(z<2 && y>1){
err=queue->enqueueNDRangeKernel(*_kernels.at(id1)._kernel, cl::NullRange/*offset*/, cl::NDRange(x,y)/*Global*/, cl::NullRange/*local*/);
}else{
err=queue->enqueueNDRangeKernel(*_kernels.at(id1)._kernel, cl::NullRange/*offset*/, cl::NDRange(x)/*Global*/, cl::NullRange/*local*/);
};
queue->finish();
if(err!=0){
std::cerr<<"kernel: "<<_kernels.at(id1)._kernel->getInfo<CL_KERNEL_FUNCTION_NAME>();
std::cerr<<" ,Args: "<<_kernels.at(id1)._kernel->getInfo<CL_KERNEL_NUM_ARGS>()<<" ,Run out: "<<err<<'\n';
for(int i=0;i<_kernels.at(id1)._kernel->getInfo<CL_KERNEL_NUM_ARGS>();++i)
std::cerr<<"argument type "<<i<<":"<<_kernels.at(id1)._kernel->getArgInfo<CL_KERNEL_ARG_TYPE_NAME>(i)<<'\n';
}
return err;
}
template<typename T>
void read(size_t id1/*kernal ID*/,size_t id2/*buffer ID*/, T* out){
queue->enqueueReadBuffer(_kernels.at(id1).read(id2).first, CL_TRUE, 0, sizeof(T) * _kernels.at(id1).read(id2).second, out);
}
void clear_buf();
private:
std::vector<kernel> _kernels;
cl::Context* context = NULL;
cl::Device device;
cl::Program* program = NULL;
cl::CommandQueue* queue=NULL;
};
#endif //VIO_OPENCL_CL_CLASS_H
kernel/double_add.cl
+4
View File
@@ -0,0 +1,4 @@
__kernel void double_add(__global const double* A, __global const double* B, __global double* C) {
int index = get_global_id(0);
C[index] = A[index] + B[index];
};
kernel/fast-best.cl
+72
View File
@@ -0,0 +1,72 @@
// Copyright (C) 2011 Dmitri Nikulin
// Copyright (C) 2011 Monash University
//
// Permission is hereby granted, free of charge, to any person
// obtaining a copy of this software and associated documentation
// files (the "Software"), to deal in the Software without
// restriction, including without limitation the rights to use,
// copy, modify, merge, publish, distribute, sublicense, and/or sell
// copies of the Software, and to permit persons to whom the
// Software is furnished to do so, subject to the following
// conditions:
//
// The above copyright notice and this permission notice shall be
// included in all copies or substantial portions of the Software.
//
// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
// EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES
// OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND
// NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT
// HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY,
// WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
// FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR
// OTHER DEALINGS IN THE SOFTWARE.
kernel void fast_best(
read_only image2d_t scores,
global int2 * corners,
global int2 * filtered,
global int * icorner
) {
// Prepare a suitable OpenCL image sampler.
sampler_t const sampler = CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
// Use global work item as 1D offset into corners.
int const ic = get_global_id(0);
int2 const xy = corners[ic];
// Read the candidate score.
int const p00 = read_imageui(scores, sampler, xy).x;
// Read other scores in a tight square around the candidate score.
int const p01 = read_imageui(scores, sampler, xy + (int2)(-1, -1)).x;
int const p02 = read_imageui(scores, sampler, xy + (int2)( 0, -1)).x;
int const p03 = read_imageui(scores, sampler, xy + (int2)( 1, -1)).x;
int const p04 = read_imageui(scores, sampler, xy + (int2)(-1, 0)).x;
int const p05 = read_imageui(scores, sampler, xy + (int2)( 0, 0)).x;
int const p06 = read_imageui(scores, sampler, xy + (int2)( 1, 0)).x;
int const p07 = read_imageui(scores, sampler, xy + (int2)(-1, 1)).x;
int const p08 = read_imageui(scores, sampler, xy + (int2)( 0, 1)).x;
int const p09 = read_imageui(scores, sampler, xy + (int2)( 1, 1)).x;
// Select the maximum score.
int sco = p00;
sco = max(sco, p01);
sco = max(sco, p02);
sco = max(sco, p03);
sco = max(sco, p04);
sco = max(sco, p05);
sco = max(sco, p06);
sco = max(sco, p07);
sco = max(sco, p08);
sco = max(sco, p09);
// Keep this score if it is as good as the maximum.
if (p00 >= sco) {
// Atomically append to filtered corner buffer.
filtered[atom_inc(icorner)] = xy;
}
}
kernel/fast-gray.cl
+80
View File
@@ -0,0 +1,80 @@
// Copyright (C) 2021 Majid Geravand
// Copyright (C) 2021 Gfuse
__kernel void fast_gray(
__read_only image2d_t image,
__global int2 * corners,
__global int * icorner,
int fast_count
) {
// Prepare a suitable OpenCL image sampler.
sampler_t const sampler = CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
// Use global work item as 2D image coordinates.
int const x = get_global_id(0);
int const y = get_global_id(1);
if(x>16 && x<get_image_width(image)-16 && y<get_image_height(image)-16 && y>16){
int2 const xy = (int2)(x, y);
// Read the candidate pixel.
int const p00 = read_imageui(image, sampler, xy).x;
// Read other pixels in a circle around the candidate pixel.
int const p01 = read_imageui(image, sampler, xy + (int2)( 0, 3)).x;
int const p05 = read_imageui(image, sampler, xy + (int2)( 3, 0)).x;
int const p09 = read_imageui(image, sampler, xy + (int2)( 0, -3)).x;
int const p13 = read_imageui(image, sampler, xy + (int2)(-3, 0)).x;
// Check the absolute difference of each circle pixel.
int const d01 = (abs(p01 - p00) > FAST_THRESH);
int const d05 = (abs(p05 - p00) > FAST_THRESH);
int const d09 = (abs(p09 - p00) > FAST_THRESH);
int const d13 = (abs(p13 - p00) > FAST_THRESH);
// Check if any two adjacent circle pixels have a high absolute difference.
if ((d01 && d05) ||(d05 && d09) ||(d09 && d13) ||(d13 && d01)) {
// Read other pixels in a circle around the candidate pixel.
int const p02 = read_imageui(image, sampler, xy + (int2)( 1, 3)).x;
int const p03 = read_imageui(image, sampler, xy + (int2)( 2, 2)).x;
int const p04 = read_imageui(image, sampler, xy + (int2)( 3, 1)).x;
int const p06 = read_imageui(image, sampler, xy + (int2)( 3, -1)).x;
int const p07 = read_imageui(image, sampler, xy + (int2)( 2, -2)).x;
int const p08 = read_imageui(image, sampler, xy + (int2)( 1, -3)).x;
int const p10 = read_imageui(image, sampler, xy + (int2)(-1, -3)).x;
int const p11 = read_imageui(image, sampler, xy + (int2)(-2, -2)).x;
int const p12 = read_imageui(image, sampler, xy + (int2)(-3, -1)).x;
int const p14 = read_imageui(image, sampler, xy + (int2)(-3, 1)).x;
int const p15 = read_imageui(image, sampler, xy + (int2)(-2, 2)).x;
int const p16 = read_imageui(image, sampler, xy + (int2)(-1, 3)).x;
// Select the maximum score.
int sco = p00;
sco = max(sco, p01);
sco = max(sco, p02);
sco = max(sco, p03);
sco = max(sco, p04);
sco = max(sco, p05);
sco = max(sco, p06);
sco = max(sco, p07);
sco = max(sco, p08);
sco = max(sco, p09);
sco = max(sco, p10);
sco = max(sco, p11);
sco = max(sco, p12);
sco = max(sco, p13);
sco = max(sco, p14);
sco = max(sco, p15);
// Keep this score if it is as good as the maximum.
if (p00 >= sco) {
// Atomically append to corner buffer.
int const icorn = atom_inc(icorner);
if (icorn < fast_count)
corners[icorn] = xy;
}
}
}
}
kernel/int_add.cl
+4
View File
@@ -0,0 +1,4 @@
__kernel void int_add(__global const int* A, __global const int* B, __global int* C) {
int index = get_global_id(0);
C[index] = A[index] + B[index];
};
kernel/is_corner_10_greater.cl
+372
View File
@@ -0,0 +1,372 @@
__kernel void is_corner_10_greater(__global const unsigned char* img,__global const int* x,__global const int* y, int2 image_stride, int wight, short barrier,__global unsigned char* condition) {
/*
int x_id=x[get_global_id(0)];
int y_id=y[get_global_id(1)];
printf("%d, %d \n",x_id,y_id);
int px=y_id*image_stride+x_id;
int w3 = 3*image_stride;//3*image_stride
/*
int t = (img[px] << 8) | img[px+1] - barrier;
if (!img[px-1-w3]>t) {
if (!img[px+3+image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+2+2*image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-1+w3]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+1+w3]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+w3]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-2+2*image_stride]>t) {
if (!img[px-w3]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+3]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+1-w3]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+2-2*image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+3-image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
condition[y_id*wight+x_id]=255;
}
if (!img[px-3+image_stride]>t) {
if (!img[px+3]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+1-w3]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+2-2*image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+3-image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
condition[y_id*wight+x_id]=255;
}
if (!img[px+3]>t) {
if (!img[px-3]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-3-image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-2-2*image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
condition[y_id*wight+x_id]=255;
}
if (!img[px+3-image_stride]>t) {
if (!img[px-3]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-3-image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
condition[y_id*wight+x_id]=255;
}
if (!img[px-3]>t) {
if (!img[px+2-2*image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
condition[y_id*wight+x_id]=255;
}
condition[y_id*wight+x_id]=255;
}
if (!img[px-2-2*image_stride]>t) {
if (!img[px+3]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+3+image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+w3]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+1+w3]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+2+2*image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+3-image_stride]>t) {
if (!img[px-1+w3]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-3-image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-2+2*image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-3+image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-3]>t) {
condition[y_id*wight+x_id]=0;
}
condition[y_id*wight+x_id]=255;
}
if (!img[px+2-2*image_stride]>t) {
if (!img[px-3]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-1+w3]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-2+2*image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-3+image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
condition[y_id*wight+x_id]=255;
}
if (!img[px+1-w3]>t) {
if (!img[px-1+w3]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-2+2*image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-3+image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
condition[y_id*wight+x_id]=255;
}
if (!img[px-w3]>t) {
if (!img[px-1+w3]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-2+2*image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
condition[y_id*wight+x_id]=255;
}
condition[y_id*wight+x_id]=255;
}
if (!img[px-3-image_stride]>t) {
if (!img[px+1+w3]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+3-image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+3]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+3+image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+2+2*image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+2-2*image_stride]>t) {
if (!img[px+w3]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-3]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-1+w3]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-2+2*image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-3+image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
condition[y_id*wight+x_id]=255;
}
if (!img[px+1-w3]>t) {
if (!img[px-3+image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+w3]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-1+w3]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-2+2*image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
condition[y_id*wight+x_id]=255;
}
if (!img[px-w3]>t) {
if (!img[px+w3]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-1+w3]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-2+2*image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
condition[y_id*wight+x_id]=255;
}
condition[y_id*wight+x_id]=255;
}
if (!img[px-w3]>t) {
if (!img[px+2+2*image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+1+w3]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-2+2*image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+w3]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-1+w3]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-3+image_stride]>t) {
if (!img[px+1-w3]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+3+image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+2-2*image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+3-image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+3]>t) {
condition[y_id*wight+x_id]=0;
}
condition[y_id*wight+x_id]=255;
}
if (!img[px-3]>t) {
if (!img[px+3+image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+2-2*image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+3-image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+3]>t) {
condition[y_id*wight+x_id]=0;
}
condition[y_id*wight+x_id]=255;
}
condition[y_id*wight+x_id]=255;
}
if (!img[px-3]>t) {
if (!img[px+2+2*image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+2-2*image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+3-image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+3+image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+3]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+1-w3]>t) {
if (!img[px+1+w3]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-3+image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+w3]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-1+w3]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-2+2*image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
condition[y_id*wight+x_id]=255;
}
condition[y_id*wight+x_id]=255;
}
if (!img[px+1-w3]>t) {
if (!img[px+1+w3]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-3+image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+w3]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-1+w3]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-2+2*image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
condition[y_id*wight+x_id]=255;
}
if (!img[px+2-2*image_stride]>t) {
if (!img[px-3+image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+w3]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-1+w3]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-2+2*image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
condition[y_id*wight+x_id]=255;
}
if (!img[px-3+image_stride]>t) {
if (!img[px+3-image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+3]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+3+image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
condition[y_id*wight+x_id]=255;
}
if (!img[px-2+2*image_stride]>t) {
if (!img[px+3-image_stride]>t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+3]>t) {
condition[y_id*wight+x_id]=0;
}
condition[y_id*wight+x_id]=255;
}
if (!img[px+3-image_stride]>t) {
if (!img[px-1+w3]>t) {
condition[y_id*wight+x_id]=0;
}
condition[y_id*wight+x_id]=255;
}
*/
};
kernel/is_corner_10_less.cl
+376
View File
@@ -0,0 +1,376 @@
__kernel void is_corner_10_less(__global const uchar* img,__global const int* x,__global const int* y,int image_stride, int wight,int barrier,
__global uchar* condition) {
int x_id=x[get_global_id(0)];
int y_id=y[get_global_id(1)];
int px=y_id*image_stride+x_id;
int w3 = 3*image_stride;//3*image_stride
int t = (img[px] << 8) | img[px+1] - barrier;
if(get_global_id(0)>100){
condition[get_global_id(1)*752+get_global_id(0)]=255;
}else{
if(condition[get_global_id(1)*752+get_global_id(0)]!=100)
condition[get_global_id(1)*752+get_global_id(0)]=0;
}
/*
if (!img[px-1-w3]<t) {
if (!img[px+3+image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+2+2*image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-1+w3]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+1+w3]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+w3]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-2+2*image_stride]<t) {
if (!img[px-w3]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+3]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+1-w3]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+2-2*image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+3-image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
condition[y_id*wight+x_id]=255;
}
if (!img[px-3+image_stride]<t) {
if (!img[px+3]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+1-w3]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+2-2*image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+3-image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
condition[y_id*wight+x_id]=255;
}
if (!img[px+3]<t) {
if (!img[px-3]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-3-image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-2-2*image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
condition[y_id*wight+x_id]=255;
}
if (!img[px+3-image_stride]<t) {
if (!img[px-3]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-3-image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
condition[y_id*wight+x_id]=255;
}
if (!img[px-3]<t) {
if (!img[px+2-2*image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
condition[y_id*wight+x_id]=255;
}
condition[y_id*wight+x_id]=255;
}
if (!img[px-2-2*image_stride]<t) {
if (!img[px+3]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+3+image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+w3]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+1+w3]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+2+2*image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+3-image_stride]<t) {
if (!img[px-1+w3]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-3-image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-2+2*image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-3+image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-3]<t) {
condition[y_id*wight+x_id]=0;
}
condition[y_id*wight+x_id]=255;
}
if (!img[px+2-2*image_stride]<t) {
if (!img[px-3]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-1+w3]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-2+2*image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-3+image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
condition[y_id*wight+x_id]=255;
}
if (!img[px+1-w3]<t) {
if (!img[px-1+w3]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-2+2*image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-3+image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
condition[y_id*wight+x_id]=255;
}
if (!img[px-w3]<t) {
if (!img[px-1+w3]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-2+2*image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
condition[y_id*wight+x_id]=255;
}
condition[y_id*wight+x_id]=255;
}
if (!img[px-3-image_stride]<t) {
if (!img[px+1+w3]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+3-image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+3]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+3+image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+2+2*image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+2-2*image_stride]<t) {
if (!img[px+w3]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-3]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-1+w3]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-2+2*image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-3+image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
condition[y_id*wight+x_id]=255;
}
if (!img[px+1-w3]<t) {
if (!img[px-3+image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+w3]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-1+w3]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-2+2*image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
condition[y_id*wight+x_id]=255;
}
if (!img[px-w3]<t) {
if (!img[px+w3]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-1+w3]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-2+2*image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
condition[y_id*wight+x_id]=255;
}
condition[y_id*wight+x_id]=255;
}
if (!img[px-w3]<t) {
if (!img[px+2+2*image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+1+w3]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-2+2*image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+w3]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-1+w3]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-3+image_stride]<t) {
if (!img[px+1-w3]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+3+image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+2-2*image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+3-image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+3]<t) {
condition[y_id*wight+x_id]=0;
}
condition[y_id*wight+x_id]=255;
}
if (!img[px-3]<t) {
if (!img[px+3+image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+2-2*image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+3-image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+3]<t) {
condition[y_id*wight+x_id]=0;
}
condition[y_id*wight+x_id]=255;
}
condition[y_id*wight+x_id]=255;
}
if (!img[px-3]<t) {
if (!img[px+2+2*image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+2-2*image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+3-image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+3+image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+3]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+1-w3]<t) {
if (!img[px+1+w3]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-3+image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+w3]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-1+w3]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-2+2*image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
condition[y_id*wight+x_id]=255;
}
condition[y_id*wight+x_id]=255;
}
if (!img[px+1-w3]<t) {
if (!img[px+1+w3]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-3+image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+w3]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-1+w3]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-2+2*image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
condition[y_id*wight+x_id]=255;
}
if (!img[px+2-2*image_stride]<t) {
if (!img[px-3+image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+w3]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-1+w3]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px-2+2*image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
condition[y_id*wight+x_id]=255;
}
if (!img[px-3+image_stride]<t) {
if (!img[px+3-image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+3]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+3+image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
condition[y_id*wight+x_id]=255;
}
if (!img[px-2+2*image_stride]<t) {
if (!img[px+3-image_stride]<t) {
condition[y_id*wight+x_id]=0;
}
if (!img[px+3]<t) {
condition[y_id*wight+x_id]=0;
}
condition[y_id*wight+x_id]=255;
}
if (!img[px+3-image_stride]<t) {
if (!img[px-1+w3]<t) {
condition[y_id*wight+x_id]=0;
}
condition[y_id*wight+x_id]=255;
}
*/
};
kernel/prefast-gray.cl
+103
View File
@@ -0,0 +1,103 @@
// Copyright (C) 2011 Dmitri Nikulin
// Copyright (C) 2011 Monash University
//
// Permission is hereby granted, free of charge, to any person
// obtaining a copy of this software and associated documentation
// files (the "Software"), to deal in the Software without
// restriction, including without limitation the rights to use,
// copy, modify, merge, publish, distribute, sublicense, and/or sell
// copies of the Software, and to permit persons to whom the
// Software is furnished to do so, subject to the following
// conditions:
//
// The above copyright notice and this permission notice shall be
// included in all copies or substantial portions of the Software.
//
// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
// EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES
// OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND
// NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT
// HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY,
// WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
// FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR
// OTHER DEALINGS IN THE SOFTWARE.
// Enable OpenCL 32-bit integer atomic functions.
__kernel void prefast_gray(
__read_only image2d_t image,
__global int2 * corners,
__global int * icorner,
int fast_count
) {
// Prepare a suitable OpenCL image sampler.
sampler_t const sampler = CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
// Use global work item as 2D image coordinates.
int const x = get_global_id(0);
int const y = get_global_id(1);
if(x>16 && x<get_image_width(image)-16 && y<get_image_height(image)-16 && y>16){
int2 const xy = (int2)(x, y);
// Read the candidate pixel.
int const p00 = read_imageui(image, sampler, xy).x;
// Read other pixels in a circle around the candidate pixel.
int const p01 = read_imageui(image, sampler, xy + (int2)( 0, 3)).x;
int const p05 = read_imageui(image, sampler, xy + (int2)( 3, 0)).x;
int const p09 = read_imageui(image, sampler, xy + (int2)( 0, -3)).x;
int const p13 = read_imageui(image, sampler, xy + (int2)(-3, 0)).x;
// Check the absolute difference of each circle pixel.
int const d01 = (abs(p01 - p00) > FAST_THRESH);
int const d05 = (abs(p05 - p00) > FAST_THRESH);
int const d09 = (abs(p09 - p00) > FAST_THRESH);
int const d13 = (abs(p13 - p00) > FAST_THRESH);
// Check if any two adjacent circle pixels have a high absolute difference.
if ((d01 && d05) ||(d05 && d09) ||(d09 && d13) ||(d13 && d01)) {
// Read other pixels in a circle around the candidate pixel.
int const p02 = read_imageui(image, sampler, xy + (int2)( 1, 3)).x;
int const p03 = read_imageui(image, sampler, xy + (int2)( 2, 2)).x;
int const p04 = read_imageui(image, sampler, xy + (int2)( 3, 1)).x;
int const p06 = read_imageui(image, sampler, xy + (int2)( 3, -1)).x;
int const p07 = read_imageui(image, sampler, xy + (int2)( 2, -2)).x;
int const p08 = read_imageui(image, sampler, xy + (int2)( 1, -3)).x;
int const p10 = read_imageui(image, sampler, xy + (int2)(-1, -3)).x;
int const p11 = read_imageui(image, sampler, xy + (int2)(-2, -2)).x;
int const p12 = read_imageui(image, sampler, xy + (int2)(-3, -1)).x;
int const p14 = read_imageui(image, sampler, xy + (int2)(-3, 1)).x;
int const p15 = read_imageui(image, sampler, xy + (int2)(-2, 2)).x;
int const p16 = read_imageui(image, sampler, xy + (int2)(-1, 3)).x;
// Select the maximum score.
int sco = p00;
sco = max(sco, p01);
sco = max(sco, p02);
sco = max(sco, p03);
sco = max(sco, p04);
sco = max(sco, p05);
sco = max(sco, p06);
sco = max(sco, p07);
sco = max(sco, p08);
sco = max(sco, p09);
sco = max(sco, p10);
sco = max(sco, p11);
sco = max(sco, p12);
sco = max(sco, p13);
sco = max(sco, p14);
sco = max(sco, p15);
// Keep this score if it is as good as the maximum.
if (p00 >= sco) {
// Atomically append to corner buffer.
int const icorn = atom_inc(icorner);
if (icorn < fast_count)
corners[icorn] = xy;
}
}
}
}
main.cpp
Executable
+96
View File
@@ -0,0 +1,96 @@
#include <iostream>
#include "cl_class.h"
#include <opencv2/opencv.hpp>
#include <opencv2/core.hpp>
#include <opencv2/core/mat.hpp>
// Maximum number of corners.
#define FAST_COUNT (1 << 18)
int main() {
opencl gpu;
/*
gpu.make_kernel("int_add");
int A[] = { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9 };
int B[] = { 0, 1, 2, 0, 1, 2, 0, 1, 2, 0 };
gpu.write_buf(0,0,10,A);
gpu.write_buf(0,1,10,B);
int C[10];
gpu.run(0,10,1,1,2,10,C);
std::cout << " result: \n";
for (int i = 0; i < 10; i++){
std::cout << C[i] << " ";
}
std::cout << std::endl;
int A1[] = { 10, 1, 2, 3, 4, 5, 6, 7, 8, 9 };
int B1[] = { 0, 1, 2, 0, 1, 2, 0, 1, 2, 10};
gpu.reload_buf(0,0,A1);
gpu.reload_buf(0,1,B1);
int C1[10];
gpu.run(0,10,1,1,2,10,C1);
std::cout << " result: \n";
for (int i = 0; i < 10; i++){
std::cout << C1[i] << " ";
}
std::cout << std::endl;
gpu.make_kernel("double_add");
double A2[] = { 0.2, 1.3, 2.5, 3.6, 4.9, 5.12, 6.5, 7.6, 8.4, 9.1 };
double B2[] = { 0.0, 1.0, 2.0, 0.0, 1.0, 2.0, 0.0, 1.0, 2.0, 0.0 };
gpu.write_buf(1,0,10,A2);
gpu.write_buf(1,1,10,B2);
double C2[10];
gpu.run(1,10,1,1,2,10,C2);
std::cout << " result: \n";
for (int i = 0; i < 10; i++){
std::cout << C2[i] << " ";
}
std::cout << std::endl;
*/
gpu.make_kernel("fast_gray");
cv::Mat img = cv::imread("../frame0287.jpg", CV_8UC1);
cv::Mat imageRGBA;
cv::Mat downSampled;
//cv::imshow("test2",img);
//cv::waitKey();
cv::resize(img, downSampled, cv::Size(752, 480));
cv::cvtColor(downSampled, imageRGBA, CV_GRAY2RGBA);
gpu.write_buf(0,0,downSampled);
//cv::imshow("test2",downSampled);
//cv::waitKey();
cl_int2 corners_[740*450];
gpu.write_buf(0,1,740*450,corners_);
int icorner[1]={0};
gpu.write_buf(0,2,1,icorner);
int fast_count=FAST_COUNT;
gpu.write_buf(0,3,fast_count);
printf("\nOpenCL version\n");
double time_accumulator = 0;
for(int i=0;i<1000;++i){
double t = (double)cv::getTickCount();
gpu.reload_buf(0,0,downSampled);
int icorner[1]={0};
gpu.reload_buf(0,2,icorner);
gpu.run(0,752,480,1);
time_accumulator += ((cv::getTickCount() - t) / cv::getTickFrequency());
}
printf("OpenCL took %f ms (average over %d trials).\n", time_accumulator/((double)1000)*1000.0, 1000 );
int icorner_[1]={0};
gpu.read(0,2,icorner_);
cl_int2 corners[740*450];
gpu.read(0,1,corners);
//double t = (double)cv::getTickCount();
//cv::Mat condition=cv::Mat::ones(cv::Size(752,480),CV_8UC1);
//gpu.run(2,752,480,1,6,360960,condition.data);
int count=0;
for(int i=0;i<icorner_[0];++i){
if(corners[i].x<0 || corners[i].x>750 || corners[i].y>450 || corners[i].y<0)continue;
cv::circle(imageRGBA, cv::Point(corners[i].x,corners[i].y ), 1, cv::Scalar(0, 0,255),1);
++count;
}
cv::imshow("test3",imageRGBA);
std::cerr<<count<<'\n';
cv::waitKey();
gpu.clear_buf();
return 0;
}
readme.md
Executable
+11
View File
@@ -0,0 +1,11 @@
sudo apt-get install -y gpg-agent wget
wget -qO - https://repositories.intel.com/graphics/intel-graphics.key |
sudo apt-key add -
sudo apt-add-repository \
'deb [arch=amd64] https://repositories.intel.com/graphics/ubuntu bionic main'
sudo apt-get update
sudo apt install ocl-icd-* opencl-headers intel-level-zero-gpu level-zero intel-opencl intel-igc-opencl-devel level-zero-devel
CodeBlocks build
cmake .. -G"CodeBlocks - Unix Makefiles"
src/FastGrayStep.cc
+71
View File
@@ -0,0 +1,71 @@
// Copyright (C) 2011 Dmitri Nikulin
// Copyright (C) 2011 Monash University
//
// Permission is hereby granted, free of charge, to any person
// obtaining a copy of this software and associated documentation
// files (the "Software"), to deal in the Software without
// restriction, including without limitation the rights to use,
// copy, modify, merge, publish, distribute, sublicense, and/or sell
// copies of the Software, and to permit persons to whom the
// Software is furnished to do so, subject to the following
// conditions:
//
// The above copyright notice and this permission notice shall be
// included in all copies or substantial portions of the Software.
//
// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
// EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES
// OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND
// NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT
// HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY,
// WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
// FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR
// OTHER DEALINGS IN THE SOFTWARE.
#include "cvd-cl/steps/FastGrayStep.hh"
#include "kernels/fast-gray.hh"
namespace CVD {
namespace CL {
FastGrayStep::FastGrayStep(GrayImageState & i_image, PointListState & i_points, PointListState & o_points, cl_int threshold, cl_int ring) :
WorkerStep (i_image.worker),
i_image (i_image),
i_points (i_points),
o_points (o_points),
threshold (threshold),
ring (ring)
{
char opt[256] = {0,};
snprintf(opt, sizeof(opt) - 1, "-DFAST_THRESH=%d -DFAST_RING=%d -DFAST_COUNT=%d", int(threshold), int(ring), int(o_points.size));
worker.compile(&program, &kernel, OCL_FAST_GRAY, "fast_gray", opt);
}
FastGrayStep::~FastGrayStep() {
// Do nothing.
}
void FastGrayStep::execute() {
// Read number of input points.
size_t const np = i_points.getCount();
// Assign kernel parameters.
kernel.setArg(0, i_image.image);
kernel.setArg(1, i_points.buffer);
kernel.setArg(2, o_points.buffer);
kernel.setArg(3, o_points.count);
kernel.setArg(4, (cl_uint) np);
// Reset number of output points.
o_points.setCount(0);
// Prepare global and local size.
cl::NDRange const global(worker.padGlobal(np));
cl::NDRange const local(worker.defaultLocal);
// Queue kernel with global size set to number of input points.
worker.queue.enqueueNDRangeKernel(kernel, cl::NullRange, global, local);
}
} // namespace CL
} // namespace CVD
src/PreFastGrayStep.cc
+70
View File
@@ -0,0 +1,70 @@
// Copyright (C) 2011 Dmitri Nikulin
// Copyright (C) 2011 Monash University
//
// Permission is hereby granted, free of charge, to any person
// obtaining a copy of this software and associated documentation
// files (the "Software"), to deal in the Software without
// restriction, including without limitation the rights to use,
// copy, modify, merge, publish, distribute, sublicense, and/or sell
// copies of the Software, and to permit persons to whom the
// Software is furnished to do so, subject to the following
// conditions:
//
// The above copyright notice and this permission notice shall be
// included in all copies or substantial portions of the Software.
//
// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
// EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES
// OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND
// NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT
// HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY,
// WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
// FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR
// OTHER DEALINGS IN THE SOFTWARE.
#include "cvd-cl/steps/PreFastGrayStep.hh"
#include "kernels/prefast-gray.hh"
namespace CVD {
namespace CL {
PreFastGrayStep::PreFastGrayStep(GrayImageState & i_image, PointListState & o_points, cl_int threshold) :
WorkerStep (i_image.worker),
i_image (i_image),
o_points (o_points),
threshold (threshold)
{
char opt[256] = {0,};
snprintf(opt, sizeof(opt) - 1, "-DFAST_THRESH=%d -DFAST_COUNT=%d", int(threshold), int(o_points.size));
worker.compile(&program, &kernel, OCL_PRE_FAST_GRAY, "prefast_gray", opt);
}
PreFastGrayStep::~PreFastGrayStep() {
// Do nothing.
}
void PreFastGrayStep::execute() {
// Reset number of output points.
o_points.setCount(0);
// Assign kernel parameters.
kernel.setArg(0, i_image.image);
kernel.setArg(1, o_points.buffer);
kernel.setArg(2, o_points.count);
// Read image dimensions.
size_t const nx = i_image.nx - 16;
size_t const ny = i_image.ny - 16;
// Construct global, local and offset with a safety boundary.
// 16x16 appears to give good performance on most devices.
cl::NDRange const global (nx - 16, ny - 16);
cl::NDRange const local ( 8, 8);
cl::NDRange const offset ( 8, 8);
// Queue kernel.
worker.queue.enqueueNDRangeKernel(kernel, offset, global, local);
}
} // namespace CL
} // namespace CVD
src/cl_class.cpp
+48
View File
@@ -0,0 +1,48 @@
//
// Created by root on 4/27/21.
//
#include "cl_class.h"
opencl::opencl() {
std::vector<cl::Platform> all_platforms;
cl::Platform::get(&all_platforms);
if (all_platforms.size() == 0){
std::cerr << " No platforms found. Check OpenCL installation!\n";
exit(1);
}
//get default device of the default platform
std::vector<cl::Device> all_devices;
for(auto i:all_platforms){
std::cout << "Find platform: " << i.getInfo<CL_PLATFORM_NAME>() << "\n";
i.getDevices(CL_DEVICE_TYPE_GPU, &all_devices);
}
if (all_devices.size() == 0){
std::cout << " No devices found. Check OpenCL installation!\n";
exit(1);
}
device = all_devices[0];
std::cout << "CL_DEVICE_NAME: " << device.getInfo<CL_DEVICE_NAME>() <<'\n'
<< "CL_DEVICE_OPENCL_C_VERSION: " <<device.getInfo<CL_DEVICE_OPENCL_C_VERSION>()<<'\n'
<< "CL_DEVICE_BUILT_IN_KERNELS: " <<device.getInfo<CL_DEVICE_BUILT_IN_KERNELS>()<<'\n'
<< "CL_DEVICE_COMPILER_AVAILABLE: " <<device.getInfo<CL_DEVICE_COMPILER_AVAILABLE>()<<'\n'
<< "CL_DEVICE_LOCAL_MEM_SIZE: " <<device.getInfo<CL_DEVICE_LOCAL_MEM_SIZE>()<<'\n'
<< "CL_DEVICE_GLOBAL_MEM_SIZE: " <<device.getInfo<CL_DEVICE_GLOBAL_MEM_SIZE>()<<'\n';
context=new cl::Context({ device });
cl::Program::Sources sources;
read_cl fast("../kernel/fast-gray.cl");
sources.push_back({ fast.src_str, fast.size });
program=new cl::Program(*context, sources);
if(program->build({ device },"-DFAST_THRESH=40") !=0)
std::cout << " Error building: " << program->getBuildInfo<CL_PROGRAM_BUILD_STATUS>(device)<<'\n'
<<program->getBuildInfo<CL_PROGRAM_BUILD_LOG>(device) << '\n';
queue=new cl::CommandQueue(*context,device,0,NULL);
}
opencl::~opencl() {
}
void opencl::clear_buf() {
queue->finish();
queue->flush();
_kernels.clear();
}
src/fastcl.cc
+135
View File
@@ -0,0 +1,135 @@
// Copyright (C) 2011 Dmitri Nikulin
// Copyright (C) 2011 Monash University
//
// Permission is hereby granted, free of charge, to any person
// obtaining a copy of this software and associated documentation
// files (the "Software"), to deal in the Software without
// restriction, including without limitation the rights to use,
// copy, modify, merge, publish, distribute, sublicense, and/or sell
// copies of the Software, and to permit persons to whom the
// Software is furnished to do so, subject to the following
// conditions:
//
// The above copyright notice and this permission notice shall be
// included in all copies or substantial portions of the Software.
//
// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
// EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES
// OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND
// NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT
// HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY,
// WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
// FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR
// OTHER DEALINGS IN THE SOFTWARE.
#include "common.h"
#include <cvd-cl/steps/PreFastGrayStep.hh>
#include <cvd-cl/steps/FastGrayStep.hh>
static void clfast(cl::Device & device, CVD::Image<CVD::byte> const & image) {
// Re-interpret image pointer.
cl_uchar const * const data = reinterpret_cast<cl_uchar const *>(image.data());
int const nx = image.size().x;
int const ny = image.size().y;
CVD::CL::Worker worker(device);
// Create states.
CVD::CL::GrayImageState cl_image(worker, ny, nx, CL_MEM_READ_ONLY);
CVD::CL::PointListState cl_corners1(worker, FAST_COUNT);
CVD::CL::PointListState cl_corners2(worker, FAST_COUNT);
// Create steps.
CVD::CL::PreFastGrayStep fast1(cl_image, cl_corners1, FAST_THRESH);
CVD::CL::FastGrayStep fast2(cl_image, cl_corners1, cl_corners2, FAST_THRESH, FAST_RING);
// Write image data.
cl_image.set(data);
// Run kernels for time.
long const time1 = time(NULL);
for (int i = 0; i < REPEAT; i++)
fast1.execute();
long const time2 = time(NULL);
for (int i = 0; i < REPEAT; i++)
fast2.execute();
long const time3 = time(NULL);
// Read number of corners.
int const ncorners1 = cl_corners1.getCount();
int const ncorners2 = cl_corners2.getCount();
// Calculate microseconds per kernel.
int const us1 = (((time2 - time1) * 1000000) / REPEAT);
int const us2 = (((time3 - time2) * 1000000) / REPEAT);
// Report timing and number of corners.
std::cerr << "OpenCL" << std::endl;
std::cerr << std::setw(12) << ncorners1 << " corners 1" << std::endl;
std::cerr << std::setw(12) << ncorners2 << " corners 2" << std::endl;
std::cerr << std::setw(12) << us1 << " microseconds 1" << std::endl;
std::cerr << std::setw(12) << us2 << " microseconds 2" << std::endl;
std::cerr << std::endl;
}
void clfast(CVD::Image<CVD::byte> const & image) {
try {
// Prepare list for all OpenCL devices on all platforms.
std::vector<cl::Device> devices;
std::vector<cl::Platform> platforms;
cl::Platform::get(&platforms);
std::cerr << "Found " << platforms.size() << " OpenCL platforms" << std::endl;
for (size_t ip = 0; ip < platforms.size(); ip++) {
cl::Platform &pf = platforms.at(ip);
std::cerr
<< pf.getInfo<CL_PLATFORM_NAME >() << " ("
<< pf.getInfo<CL_PLATFORM_VENDOR >() << ", "
<< pf.getInfo<CL_PLATFORM_VERSION>() << ")"
<< std::endl;
std::vector<cl::Device> newDevices;
try {
pf.getDevices(CL_DEVICE_TYPE_ALL, &newDevices);
std::cerr << " Found " << newDevices.size() << " OpenCL devices" << std::endl;
} catch (cl::Error & err) {
std::cerr << err.what() << " (code " << err.err() << ")" << std::endl;
}
for (size_t id = 0; id < newDevices.size(); id++) {
cl::Device &dev = newDevices.at(id);
std::cerr << " " << dev.getInfo<CL_DEVICE_NAME>() << std::endl;
// Add to list of all devices.
devices.push_back(dev);
}
}
std::cerr << std::endl << std::endl;
for (size_t id = 0; id < devices.size(); id++) {
cl::Device &dev = devices.at(id);
std::cerr << "Running pipeline for \"" << dev.getInfo<CL_DEVICE_NAME>() << "\"" << std::endl;
try {
clfast(dev, image);
} catch (cl::Error & err) {
std::cerr << err.what() << " (code " << err.err() << ")" << std::endl;
}
}
} catch (cl::Error & err) {
std::cerr << err.what() << " (code " << err.err() << ")" << std::endl;
}
}
test/data/test1.png
BIN
View File
Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 354 KiB