| 515 | // there is more than one dominant orientation at a given feature location. |
| 516 | template<typename T> |
| 517 | __global__ void calcOrientation( |
| 518 | float* x_out, float* y_out, unsigned* layer_out, float* response_out, |
| 519 | float* size_out, float* ori_out, unsigned* counter, const float* x_in, |
| 520 | const float* y_in, const unsigned* layer_in, const float* response_in, |
| 521 | const float* size_in, const unsigned total_feat, |
| 522 | const CParam<T> gauss_octave, const unsigned max_feat, |
| 523 | const unsigned octave, const bool double_input) { |
| 524 | const int tid_x = threadIdx.x; |
| 525 | const int tid_y = threadIdx.y; |
| 526 | const int bsz_x = blockDim.x; |
| 527 | const int bsz_y = blockDim.y; |
| 528 | |
| 529 | const unsigned f = blockIdx.y * bsz_y + tid_y; |
| 530 | |
| 531 | const int n = ORI_HIST_BINS; |
| 532 | |
| 533 | SharedMemory<float> shared; |
| 534 | float* shrdMem = shared.getPointer(); |
| 535 | float* hist = shrdMem; |
| 536 | float* temphist = shrdMem + n * 8; |
| 537 | |
| 538 | // Initialize temporary histogram |
| 539 | for (int i = tid_x; i < ORI_HIST_BINS; i += bsz_x) |
| 540 | hist[tid_y * n + i] = 0.f; |
| 541 | __syncthreads(); |
| 542 | |
| 543 | float real_x, real_y, response, size; |
| 544 | unsigned layer; |
| 545 | |
| 546 | if (f < total_feat) { |
| 547 | // Load keypoint information |
| 548 | real_x = x_in[f]; |
| 549 | real_y = y_in[f]; |
| 550 | layer = layer_in[f]; |
| 551 | response = response_in[f]; |
| 552 | size = size_in[f]; |
| 553 | |
| 554 | const int pt_x = (int)round(real_x / (1 << octave)); |
| 555 | const int pt_y = (int)round(real_y / (1 << octave)); |
| 556 | |
| 557 | // Calculate auxiliary parameters |
| 558 | const float scl_octv = size * 0.5f / (1 << octave); |
| 559 | const int radius = (int)round(ORI_RADIUS * scl_octv); |
| 560 | const float sigma = ORI_SIG_FCTR * scl_octv; |
| 561 | const int len = (radius * 2 + 1); |
| 562 | const float exp_denom = 2.f * sigma * sigma; |
| 563 | |
| 564 | const int dim0 = gauss_octave.dims[0]; |
| 565 | const int dim1 = gauss_octave.dims[1]; |
| 566 | const int imel = dim0 * dim1; |
| 567 | |
| 568 | // Points img to correct Gaussian pyramid layer |
| 569 | const T* img_ptr = gauss_octave.ptr + layer * imel; |
| 570 | |
| 571 | // Calculate orientation histogram |
| 572 | for (int l = tid_x; l < len * len; l += bsz_x) { |
| 573 | int i = l / len - radius; |
| 574 | int j = l % len - radius; |