public/vio
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

add fast kernel

Browse Source
This commit is contained in:
gfuse
2021-05-05 19:12:01 +02:00
parent 5687fe5057
commit 5a693259d6
6 changed files with 85 additions and 117 deletions
Download Patch File Download Diff File Expand all files Collapse all files
GPU_version/svo/include/svo/cl_class.h
+42 -50
View File
@@ -10,7 +10,6 @@
#include <algorithm> #include <algorithm>
#include <fstream> #include <fstream>
#include <exception> #include <exception>
#define __CL_ENABLE_EXCEPTIONS
#include <CL/opencl.h> #include <CL/opencl.h>
#include <opencv2/opencv.hpp> #include <opencv2/opencv.hpp>
@@ -18,45 +17,40 @@ class kernel{
public: public:
kernel(cl::Program* program,std::string name){_kernel = new cl::Kernel(*program,name.c_str());}; kernel(cl::Program* program,std::string name){_kernel = new cl::Kernel(*program,name.c_str());};
~kernel(){ ~kernel(){
//for(auto i:_buffers)i.first.first.;
_buffers.clear(); _buffers.clear();
_images.clear(); _images.clear();
}; };
template<typename T> template<typename T>
int32_t write(size_t id/*buffer ID*/,T* buf,cl::CommandQueue* queue,cl::Context* context,size_t buf_size){ int32_t write(size_t id/*buffer ID*/,T* buf,cl::CommandQueue* queue,cl::Context* context,size_t buf_size){
try{ _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),
_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)); buf_size),id));
queue->enqueueWriteBuffer(_buffers.back().first.first,CL_TRUE, 0, sizeof(T) * buf_size, buf); cl_int err=queue->enqueueWriteBuffer(_buffers.back().first.first,CL_TRUE, 0, sizeof(T) * buf_size, buf);
_kernel->setArg(id,_buffers.back().first.first); if(err!=0){
return CL_SUCCESS; std::cerr<<"Error C:Kernel, F:write, L:enqueueWriteBuffer\n"<<err<<'\n';
}catch (std::exception& err){ return err;
std::cerr<<"Error C:Kernel, F:write\n"<<err.what()<<'\n';
return CL_MAP_FAILURE;
} }
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){ int32_t write(size_t id/*buffer ID*/,cv::Mat& buf,cl::Context* context){
try{ _images.push_back(std::pair<cl::Image2D,size_t>(cl::Image2D(*context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
//if(id==0){
_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), cl::ImageFormat(CL_R, CL_UNSIGNED_INT8),
buf.size().width, buf.size().width,
buf.size().height, buf.size().height,
0, 0,
reinterpret_cast<uchar*>(buf.data)),id)); reinterpret_cast<uchar*>(buf.data)),id));
/* }else{
_images.push_back(std::pair<cl::Image2D,size_t>(cl::Image2D(*context, CL_MEM_READ_WRITE,
cl::ImageFormat(CL_RGBA, CL_UNORM_INT8),
buf.size().width,
buf.size().height),id));
}*/
_kernel->setArg(id,_images.back().first); cl_int err=_kernel->setArg(id,_images.back().first);
return CL_SUCCESS; if(err!=0){
}catch (std::exception& err){ std::cerr<<"Error C:Kernel, F:write image, L:setArg\n"<<err<<'\n';
std::cerr<<"Error C:Kernel, F:write image\n"<<err.what()<<'\n'; return err;
return CL_MAP_FAILURE; }
}; return CL_SUCCESS;
}; };
template<typename T> template<typename T>
@@ -87,6 +81,12 @@ public:
} }
}; };
cl::Kernel* _kernel = NULL; 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: private:
std::vector<std::pair<std::pair<cl::Buffer,size_t>,size_t>> _buffers; std::vector<std::pair<std::pair<cl::Buffer,size_t>,size_t>> _buffers;
std::vector<std::pair<cl::Image2D,size_t>> _images; std::vector<std::pair<cl::Image2D,size_t>> _images;
@@ -132,7 +132,7 @@ public:
} }
} }
template<typename T> template<typename T>
int32_t write_buf(size_t id1/*kernal ID*/,size_t id2/*buffer ID*/,T& value){ int32_t write_buf(size_t id1/*kernal ID*/,size_t id2/*buffer ID*/,T&& value){
_kernels.at(id1)._kernel->setArg(id2,value); _kernels.at(id1)._kernel->setArg(id2,value);
}; };
template<typename T> template<typename T>
@@ -154,37 +154,29 @@ public:
return CL_MAP_FAILURE; return CL_MAP_FAILURE;
} }
} }
template<typename T> int32_t run(size_t id1/*kernal ID*/,std::size_t x=1,std::size_t y=1,std::size_t z=1) {
int32_t run(size_t id1/*kernal ID*/,std::size_t x=1,std::size_t y=1,std::size_t z=1,size_t id2=0/*last id of the kernel related to the output*/, cl_int err=0;
size_t size_out=0/*out put buffer size*/,
T* out= nullptr) {
try{
//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: ";
cl::Buffer out_buf(*context, CL_MEM_READ_WRITE, sizeof(T) * size_out);
_kernels.at(id1)._kernel->setArg(id2,out_buf);
if(z>1 && y>1){ if(z>1 && y>1){
std::cerr<<queue->enqueueNDRangeKernel(*_kernels.at(id1)._kernel, cl::NullRange/*offset*/, cl::NDRange(x,y,z)/*Global*/, cl::NullRange/*local*/)<<'\n'; 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){ }else if(z<2 && y>1){
queue->enqueueNDRangeKernel(*_kernels.at(id1)._kernel, cl::NullRange/*offset*/, cl::NDRange(x,y)/*Global*/, cl::NullRange/*local*/); err=queue->enqueueNDRangeKernel(*_kernels.at(id1)._kernel, cl::NullRange/*offset*/, cl::NDRange(x,y)/*Global*/, cl::NullRange/*local*/);
//std::cerr<<queue->enqueueNDRangeKernel(*_kernels.at(id1)._kernel, cl::NullRange/*offset*/, cl::NDRange(x,y)/*Global*/, cl::NullRange/*local*/)<<'\n';
/*std::cerr<<_kernels.at(id1)._kernel->getArgInfo<CL_KERNEL_ARG_TYPE_NAME>(0)<<'\n';
std::cerr<<_kernels.at(id1)._kernel->getArgInfo<CL_KERNEL_ARG_TYPE_NAME>(1)<<'\n';
std::cerr<<_kernels.at(id1)._kernel->getArgInfo<CL_KERNEL_ARG_TYPE_NAME>(2)<<'\n';
std::cerr<<_kernels.at(id1)._kernel->getArgInfo<CL_KERNEL_ARG_TYPE_NAME>(3)<<'\n';*/
}else{ }else{
std::cerr<<queue->enqueueNDRangeKernel(*_kernels.at(id1)._kernel, cl::NullRange/*offset*/, cl::NDRange(x)/*Global*/, cl::NullRange/*local*/)<<'\n'; err=queue->enqueueNDRangeKernel(*_kernels.at(id1)._kernel, cl::NullRange/*offset*/, cl::NDRange(x)/*Global*/, cl::NullRange/*local*/);
}; };
queue->enqueueReadBuffer(out_buf, CL_TRUE, 0, sizeof(T) * size_out, out);
//wait for kernel to finish
queue->finish(); queue->finish();
return CL_SUCCESS; if(err!=0){
} catch (std::exception& err) { std::cerr<<"kernel: "<<_kernels.at(id1)._kernel->getInfo<CL_KERNEL_FUNCTION_NAME>();
std::cerr<<"RUN kernel Error\n"<<err.what()<<'\n'; std::cerr<<" ,Args: "<<_kernels.at(id1)._kernel->getInfo<CL_KERNEL_NUM_ARGS>()<<" ,Run out: "<<err<<'\n';
return CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST; 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(); void clear_buf();
private: private:
std::vector<kernel> _kernels; std::vector<kernel> _kernels;
GPU_version/svo/include/svo/feature_detection.h
+6 -1
View File
@@ -19,6 +19,7 @@
#include <svo/global.h> #include <svo/global.h>
#include <svo/frame.h> #include <svo/frame.h>
#include <svo/cl_class.h>
namespace svo { namespace svo {
@@ -49,7 +50,10 @@ public:
const int cell_size, const int cell_size,
const int n_pyr_levels); const int n_pyr_levels);
virtual ~AbstractDetector() {}; virtual ~AbstractDetector() {
gpu_fast_->clear_buf();
gpu_fast_->make_kernel("fast_gray");
};
virtual void detect( virtual void detect(
Frame* frame, Frame* frame,
@@ -62,6 +66,7 @@ public:
/// Set grid cells of existing features as occupied /// Set grid cells of existing features as occupied
void setExistingFeatures(const Features& fts); void setExistingFeatures(const Features& fts);
opencl* gpu_fast_= nullptr;
protected: protected:
GPU_version/svo/kernel/fast-gray.cl
+2 -25
View File
@@ -1,28 +1,5 @@
// Copyright (C) 2011 Dmitri Nikulin // Copyright (C) 2021 Majid Geravand
// Copyright (C) 2011 Monash University // Copyright (C) 2021 Gfuse
//
// 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 fast_gray( __kernel void fast_gray(
__read_only image2d_t image, __read_only image2d_t image,
GPU_version/svo/src/cl_class.cpp
+7 -15
View File
@@ -25,32 +25,24 @@ opencl::opencl() {
<< "CL_DEVICE_OPENCL_C_VERSION: " <<device.getInfo<CL_DEVICE_OPENCL_C_VERSION>()<<'\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_BUILT_IN_KERNELS: " <<device.getInfo<CL_DEVICE_BUILT_IN_KERNELS>()<<'\n'
<< "CL_DEVICE_COMPILER_AVAILABLE: " <<device.getInfo<CL_DEVICE_COMPILER_AVAILABLE>()<<'\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_LOCAL_MEM_SIZE: " <<device.getInfo<CL_DEVICE_LOCAL_MEM_SIZE>()<<'\n'
<< "CL_DEVICE_GLOBAL_MEM_SIZE: " <<device.getInfo<CL_DEVICE_GLOBAL_MEM_SIZE>()<<'\n'; << "CL_DEVICE_GLOBAL_MEM_SIZE: " <<device.getInfo<CL_DEVICE_GLOBAL_MEM_SIZE>()<<'\n';
context=new cl::Context({ device }); context=new cl::Context({ device });
cl::Program::Sources sources; cl::Program::Sources sources;
read_cl int_add("../kernel/int_add.cl"); read_cl fast("../kernel/fast-gray.cl");
sources.push_back({ int_add.src_str, int_add.size }); sources.push_back({ fast.src_str, fast.size });
read_cl double_add("../kernel/double_add.cl");
sources.push_back({ double_add.src_str, double_add.size });
read_cl corner_10_less("../kernel/is_corner_10_less.cl");
sources.push_back({ corner_10_less.src_str, corner_10_less.size });
read_cl prefast("../kernel/prefast-gray.cl");
sources.push_back({ prefast.src_str, prefast.size });
//read_cl fast("../kernel/fast-gray.cl");
// sources.push_back({ fast.src_str, fast.size});
program=new cl::Program(*context, sources); program=new cl::Program(*context, sources);
if(program->build({ device },"-DFAST_THRESH=40 -DFAST_RING=15") !=0) if(program->build({ device },"-DFAST_THRESH=40") !=0)
std::cout << " Error building: " << program->getBuildInfo<CL_PROGRAM_BUILD_STATUS>(device)<<'\n' std::cout << " Error building: " << program->getBuildInfo<CL_PROGRAM_BUILD_STATUS>(device)<<'\n'
<<program->getBuildInfo<CL_PROGRAM_BUILD_LOG>(device) << '\n'; <<program->getBuildInfo<CL_PROGRAM_BUILD_LOG>(device) << '\n';
queue=new cl::CommandQueue(*context,device,0,NULL); queue=new cl::CommandQueue(*context,device,0,NULL);
} }
opencl::~opencl() { opencl::~opencl() {
queue->finish();
} }
void opencl::clear_buf() { void opencl::clear_buf() {
_kernels.clear(); queue->finish();
queue->flush(); queue->flush();
_kernels.clear();
} }
GPU_version/svo/src/feature_detection.cpp
+28 -25
View File
@@ -16,7 +16,6 @@
#include <svo/feature_detection.h> #include <svo/feature_detection.h>
#include <svo/feature.h> #include <svo/feature.h>
#include <fast/fast.h>
#include <vikit/vision.h> #include <vikit/vision.h>
namespace svo { namespace svo {
@@ -32,7 +31,16 @@ AbstractDetector::AbstractDetector(
grid_n_cols_(ceil(static_cast<double>(img_width)/cell_size_)), grid_n_cols_(ceil(static_cast<double>(img_width)/cell_size_)),
grid_n_rows_(ceil(static_cast<double>(img_height)/cell_size_)), grid_n_rows_(ceil(static_cast<double>(img_height)/cell_size_)),
grid_occupancy_(grid_n_cols_*grid_n_rows_, false) grid_occupancy_(grid_n_cols_*grid_n_rows_, false)
{} {
gpu_fast_= new opencl();
cv::Mat init=cv::Mat::zeros(cv::Size(img_width, img_height),CV_8UC1);
gpu_fast_->write_buf(0,0,init);
cl_int2 corners[img_width*img_height];
gpu_fast_->write_buf(0,1,img_width*img_height,corners);
int icorner[1]={0};
gpu_fast_->write_buf(0,2,1,icorner);
gpu_fast_->write_buf(0,3,(1 << 18));
}
void AbstractDetector::resetGrid() void AbstractDetector::resetGrid()
{ {
@@ -73,35 +81,30 @@ void FastDetector::detect(
for(int L=0; L<n_pyr_levels_; ++L) for(int L=0; L<n_pyr_levels_; ++L)
{ {
const int scale = (1<<L); const int scale = (1<<L);
vector<fast::fast_xy> fast_corners; gpu_fast_->reload_buf(0,0,img_pyr[L].data);
#if __SSE2__ int icorner[1]={0};
fast::fast_corner_detect_10_sse2( gpu_fast_->reload_buf(0,2,icorner);
(fast::fast_byte*) img_pyr[L].data, img_pyr[L].cols, gpu_fast_->run(0,img_pyr[L].cols,img_pyr[L].rows);
img_pyr[L].rows, img_pyr[L].cols, 20, fast_corners); cl_int2* fast_corners;
#elif HAVE_FAST_NEON gpu_fast_->read(0,1,fast_corners);
fast::fast_corner_detect_9_neon( int count[1]={0};
(fast::fast_byte*) img_pyr[L].data, img_pyr[L].cols, gpu_fast_->read(0,2,count);
img_pyr[L].rows, img_pyr[L].cols, 20, fast_corners); //vector<int> scores, nm_corners;
#else //fast::fast_corner_score_10((fast::fast_byte*) img_pyr[L].data, img_pyr[L].cols, fast_corners, 20, scores);
fast::fast_corner_detect_10( //fast::fast_nonmax_3x3(fast_corners, scores, nm_corners);
(fast::fast_byte*) img_pyr[L].data, img_pyr[L].cols,
img_pyr[L].rows, img_pyr[L].cols, 20, fast_corners);
#endif
vector<int> scores, nm_corners;
fast::fast_corner_score_10((fast::fast_byte*) img_pyr[L].data, img_pyr[L].cols, fast_corners, 20, scores);
fast::fast_nonmax_3x3(fast_corners, scores, nm_corners);
for(auto it=nm_corners.begin(), ite=nm_corners.end(); it!=ite; ++it) for(uint i=0;i<count[0];++i)
{ {
fast::fast_xy& xy = fast_corners.at(*it); if(fast_corners[i].x<0 || fast_corners[i].x>img_pyr[L].cols || fast_corners[i].y>img_pyr[L].rows || fast_corners[i].y<0)continue;
const int k = static_cast<int>((xy.y*scale)/cell_size_)*grid_n_cols_ const int k = static_cast<int>((fast_corners[i].y*scale)/cell_size_)*grid_n_cols_
+ static_cast<int>((xy.x*scale)/cell_size_); + static_cast<int>((fast_corners[i].x*scale)/cell_size_);
if(grid_occupancy_[k]) if(grid_occupancy_[k])
continue; continue;
const float score = vk::shiTomasiScore(img_pyr[L], xy.x, xy.y); const float score = vk::shiTomasiScore(img_pyr[L], fast_corners[i].x, fast_corners[i].y);
if(score > corners.at(k).score) if(score > corners.at(k).score)
corners.at(k) = Corner(xy.x*scale, xy.y*scale, score, L, 0.0f); corners.at(k) = Corner(fast_corners[i].x*scale, fast_corners[i].y*scale, score, L, 0.0f);
} }
} }
// Create feature for every corner that has high enough corner score // Create feature for every corner that has high enough corner score
GPU_version/svo/src/frame.cpp
-1
View File
@@ -22,7 +22,6 @@
#include <vikit/math_utils.h> #include <vikit/math_utils.h>
#include <vikit/vision.h> #include <vikit/vision.h>
#include <vikit/performance_monitor.h> #include <vikit/performance_monitor.h>
#include <fast/fast.h>
namespace svo { namespace svo {