| 684 | // of Lowe's paper. |
| 685 | template<typename T> |
| 686 | __global__ void computeDescriptor( |
| 687 | float* desc_out, const unsigned desc_len, const unsigned histsz, |
| 688 | const float* x_in, const float* y_in, const unsigned* layer_in, |
| 689 | const float* response_in, const float* size_in, const float* ori_in, |
| 690 | const unsigned total_feat, const CParam<T> gauss_octave, const int d, |
| 691 | const int n, const float scale, const int n_layers) { |
| 692 | const int tid_x = threadIdx.x; |
| 693 | const int tid_y = threadIdx.y; |
| 694 | const int bsz_x = blockDim.x; |
| 695 | const int bsz_y = blockDim.y; |
| 696 | |
| 697 | const int f = blockIdx.y * bsz_y + tid_y; |
| 698 | |
| 699 | SharedMemory<float> shared; |
| 700 | float* shrdMem = shared.getPointer(); |
| 701 | float* desc = shrdMem; |
| 702 | float* accum = shrdMem + desc_len * histsz; |
| 703 | |
| 704 | for (int i = tid_x; i < desc_len * histsz; i += bsz_x) |
| 705 | desc[tid_y * desc_len + i] = 0.f; |
| 706 | __syncthreads(); |
| 707 | |
| 708 | if (f < total_feat) { |
| 709 | const unsigned layer = layer_in[f]; |
| 710 | float ori = (360.f - ori_in[f]) * PI_VAL / 180.f; |
| 711 | ori = (ori > PI_VAL) ? ori - PI_VAL * 2 : ori; |
| 712 | const float size = size_in[f]; |
| 713 | const int fx = round(x_in[f] * scale); |
| 714 | const int fy = round(y_in[f] * scale); |
| 715 | |
| 716 | const int dim0 = gauss_octave.dims[0]; |
| 717 | const int dim1 = gauss_octave.dims[1]; |
| 718 | const int imel = dim0 * dim1; |
| 719 | |
| 720 | // Points img to correct Gaussian pyramid layer |
| 721 | const T* img_ptr = gauss_octave.ptr + layer * imel; |
| 722 | |
| 723 | float cos_t = cosf(ori); |
| 724 | float sin_t = sinf(ori); |
| 725 | float bins_per_rad = n / (PI_VAL * 2.f); |
| 726 | float exp_denom = d * d * 0.5f; |
| 727 | float hist_width = DESCR_SCL_FCTR * size * scale * 0.5f; |
| 728 | int radius = hist_width * sqrtf(2.f) * (d + 1.f) * 0.5f + 0.5f; |
| 729 | |
| 730 | int len = radius * 2 + 1; |
| 731 | const int hist_off = (tid_x % histsz) * desc_len; |
| 732 | |
| 733 | // Calculate orientation histogram |
| 734 | for (int l = tid_x; l < len * len; l += bsz_x) { |
| 735 | int i = l / len - radius; |
| 736 | int j = l % len - radius; |
| 737 | |
| 738 | int y = fy + i; |
| 739 | int x = fx + j; |
| 740 | |
| 741 | float x_rot = (j * cos_t - i * sin_t) / hist_width; |
| 742 | float y_rot = (j * sin_t + i * cos_t) / hist_width; |
| 743 | float xbin = x_rot + d / 2 - 0.5f; |
nothing calls this directly
no test coverage detected