基于cuda的图像拼接，opencv

#include "akaze.h" #include "akazed.h" #include "fed.h" #include <memory> #include <cmath> //#define DEBUG_SHOW #ifdef DEBUG_SHOW #include <opencv2/core.hpp> #endif // DEBUG_SHOW #ifndef MIN # define MIN(a,b) ((a) > (b) ? (b) : (a)) #endif #ifndef MAX # define MAX(a,b) ((a) < (b) ? (b) : (a)) #endif namespace akaze { void initAkazeData(AkazeData& data, const int max_pts, const bool host, const bool dev) { data.num_pts = 0; data.max_pts = max_pts; const size_t size = sizeof(AkazePoint) * max_pts; data.h_data = host ? (AkazePoint*)malloc(size) : NULL; data.d_data = NULL; if (dev) { CHECK(cudaMalloc((void**)&data.d_data, size)); } } void freeAkazeData(AkazeData& data) { if (data.d_data != NULL) { CHECK(cudaFree(data.d_data)); } if (data.h_data != NULL) { free(data.h_data); } data.num_pts = 0; data.max_pts = 0; } void cuMatch(AkazeData& result1, AkazeData& result2) { hMatch(result1, result2); if (result1.h_data) { int* h_ptr = &result1.h_data[0].match; int* d_ptr = &result1.d_data[0].match; CHECK(cudaMemcpy2D(h_ptr, sizeof(AkazePoint), d_ptr, sizeof(AkazePoint), 4 * sizeof(float), result1.num_pts, cudaMemcpyDeviceToHost)); } } Akazer::Akazer() { } Akazer::~Akazer() { CHECK(cudaFree(omem)); } void Akazer::init(int3 whp0, int _noctaves, int _max_scale, float _per, float _kcontrast, float _soffset, bool _reordering, float _derivative_factor, float _dthreshold, int _diffusivity, int _descriptor_pattern_size) { whp.x = whp0.x; whp.y = whp0.y; whp.z = whp0.z; noctaves = _noctaves; max_scale = _max_scale; per = _per; kcontrast = _kcontrast; soffset = _soffset; reordering = _reordering; derivative_factor = _derivative_factor; dthreshold = _dthreshold; diffusivity = DiffusivityType(_diffusivity); descriptor_pattern_size = _descriptor_pattern_size; setCompareIndices(); } void Akazer::detectAndCompute(float* image, AkazeData& result, int3 whp0, const bool desc) { // Allocate memory std::unique_ptr<int> oparams(new int[noctaves * 5 + 1]); int* osizes = oparams.get(); int* offsets = osizes + noctaves; int3* owhps = (int3*)(offsets + noctaves + 1); float* tmem = NULL; const bool reused = whp0.x == whp.x && whp0.y == whp.y; if (reused) { this->allocMemory((void**)&omem, whp0, owhps, osizes, offsets, reused); tmem = omem; } else { this->allocMemory((void**)&tmem, whp0, owhps, osizes, offsets, reused); } // Detect keypoints this->detect(result, tmem, image, owhps, osizes, offsets); // Compute descriptors if (desc) { // Compute orientations hCalcOrient(result, tmem, noctaves, max_scale); // Compute descriptors hDescribe(result, tmem, noctaves, max_scale, descriptor_pattern_size); } // Copy point data to host if (result.h_data != NULL) { float* h_ptr = &result.h_data[0].x; float* d_ptr = &result.d_data[0].x; CHECK(cudaMemcpy2D(h_ptr, sizeof(AkazePoint), d_ptr, sizeof(AkazePoint), (desc ? FLEN * sizeof(unsigned char) : 0) + 6 * sizeof(float), result.num_pts, cudaMemcpyDeviceToHost)); } // Post-processing if (reused) { CHECK(cudaMemset(omem, 0, total_osize)); } else { CHECK(cudaFree(tmem)); } } void Akazer::fastDetectAndCompute(unsigned char* image, AkazeData& result, int3 whp0, const bool desc) { // Allocate memory std::unique_ptr<int> oparams(new int[noctaves * 5 + 1]); int* osizes = oparams.get(); int* offsets = osizes + noctaves; int3* owhps = (int3*)(offsets + noctaves + 1); void* tmem = NULL; const bool reused = whp0.x == whp.x && whp0.y == whp.y; if (reused) { this->allocMemory((void**)&omem, whp0, owhps, osizes, offsets, reused); tmem = omem; } else { this->allocMemory((void**)&tmem, whp0, owhps, osizes, offsets, reused); } // Detect keypoints this->fastDetect(result, tmem, image, owhps, osizes, offsets); // Compute descriptors if (desc) { // Compute orientations fastakaze::hCalcOrient(result, tmem, noctaves, max_scale); // Compute descriptors fastakaze::hDescribe(result, tmem, noctaves, max_scale, descriptor_pattern_size); } // Copy point data to host if (result.h_data != NULL) { float* h_ptr = &result.h_data[0].x; float* d_ptr = &result.d_data[0].x; CHECK(cudaMemcpy2D(h_ptr, sizeof(AkazePoint), d_ptr, sizeof(AkazePoint), (desc ? FLEN * sizeof(unsigned char) : 0) + 6 * sizeof(float), result.num_pts, cudaMemcpyDeviceToHost)); } // Post-processing if (reused) { CHECK(cudaMemset(omem, 0, total_osize)); } else { CHECK(cudaFree(tmem)); } } void Akazer::allocMemory(void** addr, int3& whp0, int3* owhps, int* osizes, int* offsets, const bool reused) { // Compute sizes owhps[0] = whp0; osizes[0] = whp0.y * whp0.z; offsets[0] = 3 * osizes[0]; // Record response, scale and octave map for NMS offsets[1] = offsets[0] + osizes[0] * max_scale * 4; // 4: Limg, Lsmooth, Lx, Ly for (int i = 0, j = 1, k = 2; j < noctaves; i++, j++, k++) { owhps[j].x = (owhps[i].x >> 1); owhps[j].y = (owhps[i].y >> 1); if (owhps[j].x < 80 || owhps[j].y < 80) { noctaves = j; break; } owhps[j].z = iAlignUp(owhps[j].x, 128); osizes[j] = owhps[j].y * owhps[j].z; offsets[k] = offsets[j] + osizes[j] * max_scale * 4; } // Allocate memory for images in octave if ((reused && !omem) || !reused) { CHECK(cudaMalloc(addr, offsets[noctaves] * sizeof(float))); } if (reused) { total_osize = offsets[noctaves] * sizeof(float); } //return offsets[noctaves]; } void Akazer::detect(AkazeData& result, float* tmem, float* image, int3* owhps, int* osizes, int* offsets) { // Get address of point counter unsigned int* d_point_counter_addr; getPointCounter((void**)&d_point_counter_addr); CHECK(cudaMemset(d_point_counter_addr, 0, sizeof(unsigned int))); setMaxNumPoints(result.max_pts); int w, h, p, msz, ms_msz, mstep; float* response_map = tmem; float* size_map = tmem + osizes[0]; int* layer_map = (int*)(size_map + osizes[0]); size_t nbytes = osizes[0] * sizeof(float); float minv = 1e-6f; int* iminv = (int*)&minv; CHECK(cudaMemset(layer_map, -1, nbytes)); CHECK(cudaMemset(response_map, *iminv, nbytes)); CHECK(cudaMemset(size_map, *iminv, nbytes)); float* oldnld = NULL; float* nldimg = NULL; float* smooth = NULL; float* flow = NULL; float* temp = NULL; float* dx = NULL; float* dy = NULL; float tmax = 0.25f; float esigma = soffset; float last_etime = 0.5 * soffset * soffset; float curr_etime = 0; float ttime = 0; int naux = 0; int oratio = 1; int sigma_size = 0; float smax = 1.0f; if (FEATURE_TYPE == 0 || FEATURE_TYPE == 1 || FEATURE_TYPE == 4 || FEATURE_TYPE == 5) { smax = 10.0 * sqrtf(2.0f); } else if (FEATURE_TYPE == 2 || FEATURE_TYPE == 3) { smax = 12.0 * sqrtf(2.0f); } std::unique_ptr<float> exptr(new float[max_scale * 2]); float* borders = exptr.get(); float* sizes = borders + max_scale; float psz = 10000; int neigh = 0; //float* threshs = borders + max_scale; #ifdef DEBUG_SHOW cv::Mat nldshow, detshow, dxshow, dyshow; cv::Mat response_show(owhps[0].y, owhps[0].x, CV_32FC1); cv::Mat size_show(owhps[0].y, owhps[0].x, CV_32FC1); cv::Mat layer_show(owhps[0].y, owhps[0].x, CV_32SC1);; #endif // DEBUG_SHOW for (int i = 0; i < noctaves; i++) { w = owhps[i].x; h = owhps[i].y; p = owhps[i].z; msz = osizes[i]; ms_msz = msz * max_scale; #ifdef DEBUG_SHOW nldshow.create(h, w, CV_32FC1); detshow.create(h, w, CV_32FC1); dxshow.create(h, w, CV_32FC1); dyshow.create(h, w, CV_32FC1); #endif // DEBUG_SHOW nldimg = tmem + offsets[i]; smooth = nldimg + ms_msz; flow = smooth + ms_msz; temp = flow + ms_msz; dx = flow; dy = temp; // Create nonlinear space for current octave layer for (int j = 0; j < max_scale; j++) { if (j == 0 && i == 0) {