| 273 | |
| 274 | template<typename T, typename convAccT> |
| 275 | void orb(unsigned* out_feat, float** d_x, float** d_y, float** d_score, |
| 276 | float** d_ori, float** d_size, unsigned** d_desc, |
| 277 | vector<unsigned>& feat_pyr, vector<float*>& d_x_pyr, |
| 278 | vector<float*>& d_y_pyr, vector<unsigned>& lvl_best, |
| 279 | vector<float>& lvl_scl, vector<Array<T>>& img_pyr, |
| 280 | const float fast_thr, const unsigned max_feat, const float scl_fctr, |
| 281 | const unsigned levels, const bool blur_img, |
| 282 | const LookupTable1D<int>& luTable) { |
| 283 | UNUSED(fast_thr); |
| 284 | UNUSED(max_feat); |
| 285 | UNUSED(scl_fctr); |
| 286 | UNUSED(levels); |
| 287 | unsigned patch_size = REF_PAT_SIZE; |
| 288 | |
| 289 | unsigned max_levels = feat_pyr.size(); |
| 290 | |
| 291 | // In future implementations, the user will be capable of passing his |
| 292 | // distribution instead of using the reference one |
| 293 | // CUDA_CHECK(cudaMemcpyToSymbolAsync(d_ref_pat, h_ref_pat, 256 * 4 * |
| 294 | // sizeof(int), 0, |
| 295 | // cudaMemcpyHostToDevice, getActiveStream())); |
| 296 | |
| 297 | vector<float*> d_score_pyr(max_levels); |
| 298 | vector<float*> d_ori_pyr(max_levels); |
| 299 | vector<float*> d_size_pyr(max_levels); |
| 300 | vector<unsigned*> d_desc_pyr(max_levels); |
| 301 | vector<unsigned*> d_idx_pyr(max_levels); |
| 302 | |
| 303 | unsigned total_feat = 0; |
| 304 | |
| 305 | // Calculate a separable Gaussian kernel |
| 306 | Array<convAccT> gauss_filter = createEmptyArray<convAccT>(dim4()); |
| 307 | if (blur_img) { |
| 308 | unsigned gauss_len = 9; |
| 309 | vector<convAccT> h_gauss(gauss_len); |
| 310 | gaussian1D(h_gauss.data(), gauss_len, 2.f); |
| 311 | dim4 gauss_dim(gauss_len); |
| 312 | gauss_filter = createHostDataArray<convAccT>(gauss_dim, h_gauss.data()); |
| 313 | CUDA_CHECK(cudaMemcpyAsync(gauss_filter.get(), h_gauss.data(), |
| 314 | h_gauss.size() * sizeof(convAccT), |
| 315 | cudaMemcpyHostToDevice, getActiveStream())); |
| 316 | CUDA_CHECK(cudaStreamSynchronize(cuda::getActiveStream())); |
| 317 | } |
| 318 | |
| 319 | for (int i = 0; i < (int)max_levels; i++) { |
| 320 | if (feat_pyr[i] == 0 || lvl_best[i] == 0) { continue; } |
| 321 | |
| 322 | // auto d_score_harris = memAlloc<float>(feat_pyr[i]); |
| 323 | dim4 score_dim(feat_pyr[i]); |
| 324 | Array<float> d_score_harris = |
| 325 | createEmptyArray<float>(score_dim); // harris_sorted |
| 326 | |
| 327 | // Calculate Harris responses |
| 328 | // Good block_size >= 7 (must be an odd number) |
| 329 | dim3 threads(THREADS_X, THREADS_Y); |
| 330 | dim3 blocks(divup(feat_pyr[i], threads.x), 1); |
| 331 | CUDA_LAUNCH((harris_response<T, false>), blocks, threads, |
| 332 | d_score_harris.get(), NULL, d_x_pyr[i], d_y_pyr[i], NULL, |
nothing calls this directly
no test coverage detected