| 84 | |
| 85 | template<typename T, bool use_scl> |
| 86 | __global__ void harris_response(float* score_out, float* size_out, |
| 87 | const float* x_in, const float* y_in, |
| 88 | const float* scl_in, const unsigned total_feat, |
| 89 | CParam<T> image, const unsigned block_size, |
| 90 | const float k_thr, const unsigned patch_size) { |
| 91 | unsigned f = blockDim.x * blockIdx.x + threadIdx.x; |
| 92 | |
| 93 | float ixx = 0.f, iyy = 0.f, ixy = 0.f; |
| 94 | float size = 0.f; |
| 95 | |
| 96 | if (f < total_feat) { |
| 97 | unsigned x, y; |
| 98 | float scl = 1.f; |
| 99 | if (use_scl) { |
| 100 | // Update x and y coordinates according to scale |
| 101 | scl = scl_in[f]; |
| 102 | x = (unsigned)round(x_in[f] * scl); |
| 103 | y = (unsigned)round(y_in[f] * scl); |
| 104 | } else { |
| 105 | x = (unsigned)round(x_in[f]); |
| 106 | y = (unsigned)round(y_in[f]); |
| 107 | } |
| 108 | |
| 109 | // Round feature size to nearest odd integer |
| 110 | size = 2.f * floor((patch_size * scl) / 2.f) + 1.f; |
| 111 | |
| 112 | // Avoid keeping features that might be too wide and might not fit on |
| 113 | // the image, sqrt(2.f) is the radius when angle is 45 degrees and |
| 114 | // represents widest case possible |
| 115 | unsigned patch_r = ceil(size * sqrt(2.f) / 2.f); |
| 116 | if (x < patch_r || y < patch_r || x >= image.dims[1] - patch_r || |
| 117 | y >= image.dims[0] - patch_r) |
| 118 | return; |
| 119 | |
| 120 | unsigned r = block_size / 2; |
| 121 | |
| 122 | unsigned block_size_sq = block_size * block_size; |
| 123 | for (unsigned k = threadIdx.y; k < block_size_sq; k += blockDim.y) { |
| 124 | int i = k / block_size - r; |
| 125 | int j = k % block_size - r; |
| 126 | |
| 127 | // Calculate local x and y derivatives |
| 128 | float ix = image.ptr[(x + i + 1) * image.dims[0] + y + j] - |
| 129 | image.ptr[(x + i - 1) * image.dims[0] + y + j]; |
| 130 | float iy = image.ptr[(x + i) * image.dims[0] + y + j + 1] - |
| 131 | image.ptr[(x + i) * image.dims[0] + y + j - 1]; |
| 132 | |
| 133 | // Accumulate second order derivatives |
| 134 | ixx += ix * ix; |
| 135 | iyy += iy * iy; |
| 136 | ixy += ix * iy; |
| 137 | } |
| 138 | } |
| 139 | __syncthreads(); |
| 140 | |
| 141 | ixx = block_reduce_sum(ixx); |
| 142 | iyy = block_reduce_sum(iyy); |
| 143 | ixy = block_reduce_sum(ixy); |
nothing calls this directly
no test coverage detected