diff --git a/GPU_version/svo/include/svo/cl_class.h b/GPU_version/svo/include/svo/cl_class.h index c0b378a..ff4cd39 100644 --- a/GPU_version/svo/include/svo/cl_class.h +++ b/GPU_version/svo/include/svo/cl_class.h @@ -10,7 +10,6 @@ #include #include #include -#define __CL_ENABLE_EXCEPTIONS #include #include @@ -18,45 +17,40 @@ class kernel{ public: kernel(cl::Program* program,std::string name){_kernel = new cl::Kernel(*program,name.c_str());}; ~kernel(){ - //for(auto i:_buffers)i.first.first.; _buffers.clear(); _images.clear(); }; template 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,size_t>(std::pair(cl::Buffer(*context, CL_MEM_READ_WRITE, sizeof(T) * buf_size), + _buffers.push_back(std::pair,size_t>(std::pair(cl::Buffer(*context, CL_MEM_READ_WRITE, sizeof(T) * buf_size), buf_size),id)); - queue->enqueueWriteBuffer(_buffers.back().first.first,CL_TRUE, 0, sizeof(T) * buf_size, buf); - _kernel->setArg(id,_buffers.back().first.first); - return CL_SUCCESS; - }catch (std::exception& err){ - std::cerr<<"Error C:Kernel, F:write\n"<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"<setArg(id,_buffers.back().first.first); + if(err!=0){ + std::cerr<<"Error C:Kernel, F:write, L3:setArg\n"<(cl::Image2D(*context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, + _images.push_back(std::pair(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(buf.data)),id)); - /* }else{ - _images.push_back(std::pair(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); - return CL_SUCCESS; - }catch (std::exception& err){ - std::cerr<<"Error C:Kernel, F:write image\n"<setArg(id,_images.back().first); + if(err!=0){ + std::cerr<<"Error C:Kernel, F:write image, L:setArg\n"< @@ -87,6 +81,12 @@ public: } }; cl::Kernel* _kernel = NULL; + std::pair 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,size_t>> _buffers; std::vector> _images; @@ -132,7 +132,7 @@ public: } } template - 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); }; template @@ -154,37 +154,29 @@ public: return CL_MAP_FAILURE; } } - template - 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*/, - size_t size_out=0/*out put buffer size*/, - T* out= nullptr) { - try{ - //std::cerr<<"kernel: "<<_kernels.at(id1)._kernel->getInfo(); - //std::cerr<<" ,Args: "<<_kernels.at(id1)._kernel->getInfo()<<" ,Run out: "; - cl::Buffer out_buf(*context, CL_MEM_READ_WRITE, sizeof(T) * size_out); - _kernels.at(id1)._kernel->setArg(id2,out_buf); + 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){ - std::cerr<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){ - queue->enqueueNDRangeKernel(*_kernels.at(id1)._kernel, cl::NullRange/*offset*/, cl::NDRange(x,y)/*Global*/, cl::NullRange/*local*/); - //std::cerr<enqueueNDRangeKernel(*_kernels.at(id1)._kernel, cl::NullRange/*offset*/, cl::NDRange(x,y)/*Global*/, cl::NullRange/*local*/)<<'\n'; - /*std::cerr<<_kernels.at(id1)._kernel->getArgInfo(0)<<'\n'; - std::cerr<<_kernels.at(id1)._kernel->getArgInfo(1)<<'\n'; - std::cerr<<_kernels.at(id1)._kernel->getArgInfo(2)<<'\n'; - std::cerr<<_kernels.at(id1)._kernel->getArgInfo(3)<<'\n';*/ + err=queue->enqueueNDRangeKernel(*_kernels.at(id1)._kernel, cl::NullRange/*offset*/, cl::NDRange(x,y)/*Global*/, cl::NullRange/*local*/); }else{ - std::cerr<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(); - return CL_SUCCESS; - } catch (std::exception& err) { - std::cerr<<"RUN kernel Error\n"<getInfo(); + std::cerr<<" ,Args: "<<_kernels.at(id1)._kernel->getInfo()<<" ,Run out: "<getInfo();++i) + std::cerr<<"argument type "<getArgInfo(i)<<'\n'; + } + return err; } + template + 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 _kernels; diff --git a/GPU_version/svo/include/svo/feature_detection.h b/GPU_version/svo/include/svo/feature_detection.h index d0990a8..6267d20 100644 --- a/GPU_version/svo/include/svo/feature_detection.h +++ b/GPU_version/svo/include/svo/feature_detection.h @@ -19,6 +19,7 @@ #include #include +#include namespace svo { @@ -49,7 +50,10 @@ public: const int cell_size, const int n_pyr_levels); - virtual ~AbstractDetector() {}; + virtual ~AbstractDetector() { + gpu_fast_->clear_buf(); + gpu_fast_->make_kernel("fast_gray"); + }; virtual void detect( Frame* frame, @@ -62,6 +66,7 @@ public: /// Set grid cells of existing features as occupied void setExistingFeatures(const Features& fts); + opencl* gpu_fast_= nullptr; protected: diff --git a/GPU_version/svo/kernel/fast-gray.cl b/GPU_version/svo/kernel/fast-gray.cl index e9ba8f4..ec36066 100644 --- a/GPU_version/svo/kernel/fast-gray.cl +++ b/GPU_version/svo/kernel/fast-gray.cl @@ -1,28 +1,5 @@ -// 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. +// Copyright (C) 2021 Majid Geravand +// Copyright (C) 2021 Gfuse __kernel void fast_gray( __read_only image2d_t image, diff --git a/GPU_version/svo/src/cl_class.cpp b/GPU_version/svo/src/cl_class.cpp index 59ba287..d32f140 100644 --- a/GPU_version/svo/src/cl_class.cpp +++ b/GPU_version/svo/src/cl_class.cpp @@ -25,32 +25,24 @@ opencl::opencl() { << "CL_DEVICE_OPENCL_C_VERSION: " <()<<'\n' << "CL_DEVICE_BUILT_IN_KERNELS: " <()<<'\n' << "CL_DEVICE_COMPILER_AVAILABLE: " <()<<'\n' - << "CL_DEVICE_LOCAL_MEM_SIZE: " <()<<'\n' - << "CL_DEVICE_GLOBAL_MEM_SIZE: " <()<<'\n'; + << "CL_DEVICE_LOCAL_MEM_SIZE: " <()<<'\n' + << "CL_DEVICE_GLOBAL_MEM_SIZE: " <()<<'\n'; context=new cl::Context({ device }); cl::Program::Sources sources; - read_cl int_add("../kernel/int_add.cl"); - sources.push_back({ int_add.src_str, int_add.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}); + 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 -DFAST_RING=15") !=0) + if(program->build({ device },"-DFAST_THRESH=40") !=0) std::cout << " Error building: " << program->getBuildInfo(device)<<'\n' <getBuildInfo(device) << '\n'; queue=new cl::CommandQueue(*context,device,0,NULL); } opencl::~opencl() { - queue->finish(); } void opencl::clear_buf() { - _kernels.clear(); + queue->finish(); queue->flush(); + _kernels.clear(); } diff --git a/GPU_version/svo/src/feature_detection.cpp b/GPU_version/svo/src/feature_detection.cpp index 98183f8..8aaac0c 100644 --- a/GPU_version/svo/src/feature_detection.cpp +++ b/GPU_version/svo/src/feature_detection.cpp @@ -16,7 +16,6 @@ #include #include -#include #include namespace svo { @@ -32,7 +31,16 @@ AbstractDetector::AbstractDetector( grid_n_cols_(ceil(static_cast(img_width)/cell_size_)), grid_n_rows_(ceil(static_cast(img_height)/cell_size_)), 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() { @@ -73,35 +81,30 @@ void FastDetector::detect( for(int L=0; L fast_corners; -#if __SSE2__ - fast::fast_corner_detect_10_sse2( - (fast::fast_byte*) img_pyr[L].data, img_pyr[L].cols, - img_pyr[L].rows, img_pyr[L].cols, 20, fast_corners); -#elif HAVE_FAST_NEON - fast::fast_corner_detect_9_neon( - (fast::fast_byte*) img_pyr[L].data, img_pyr[L].cols, - img_pyr[L].rows, img_pyr[L].cols, 20, fast_corners); -#else - fast::fast_corner_detect_10( - (fast::fast_byte*) img_pyr[L].data, img_pyr[L].cols, - img_pyr[L].rows, img_pyr[L].cols, 20, fast_corners); -#endif - vector 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); + gpu_fast_->reload_buf(0,0,img_pyr[L].data); + int icorner[1]={0}; + gpu_fast_->reload_buf(0,2,icorner); + gpu_fast_->run(0,img_pyr[L].cols,img_pyr[L].rows); + cl_int2* fast_corners; + gpu_fast_->read(0,1,fast_corners); + int count[1]={0}; + gpu_fast_->read(0,2,count); + //vector 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((xy.y*scale)/cell_size_)*grid_n_cols_ - + static_cast((xy.x*scale)/cell_size_); + 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((fast_corners[i].y*scale)/cell_size_)*grid_n_cols_ + + static_cast((fast_corners[i].x*scale)/cell_size_); if(grid_occupancy_[k]) 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) - 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 diff --git a/GPU_version/svo/src/frame.cpp b/GPU_version/svo/src/frame.cpp index 44ddab6..440c32d 100644 --- a/GPU_version/svo/src/frame.cpp +++ b/GPU_version/svo/src/frame.cpp @@ -22,7 +22,6 @@ #include #include #include -#include namespace svo {