| 96 | |
| 97 | template<typename T> |
| 98 | __global__ void non_maximal(float* x_out, float* y_out, float* resp_out, |
| 99 | unsigned* count, const unsigned idim0, |
| 100 | const unsigned idim1, const T* resp_in, |
| 101 | const float min_resp, const unsigned border_len, |
| 102 | const unsigned max_corners) { |
| 103 | // Responses on the border don't have 8-neighbors to compare, discard them |
| 104 | const unsigned r = border_len + 1; |
| 105 | |
| 106 | const unsigned x = blockDim.x * blockIdx.x + threadIdx.x + r; |
| 107 | const unsigned y = blockDim.y * blockIdx.y + threadIdx.y + r; |
| 108 | |
| 109 | if (x < idim1 - r && y < idim0 - r) { |
| 110 | const T v = resp_in[x * idim0 + y]; |
| 111 | |
| 112 | // Find maximum neighborhood response |
| 113 | T max_v; |
| 114 | max_v = max_val(resp_in[(x - 1) * idim0 + y - 1], |
| 115 | resp_in[x * idim0 + y - 1]); |
| 116 | max_v = max_val(max_v, resp_in[(x + 1) * idim0 + y - 1]); |
| 117 | max_v = max_val(max_v, resp_in[(x - 1) * idim0 + y]); |
| 118 | max_v = max_val(max_v, resp_in[(x + 1) * idim0 + y]); |
| 119 | max_v = max_val(max_v, resp_in[(x - 1) * idim0 + y + 1]); |
| 120 | max_v = max_val(max_v, resp_in[(x)*idim0 + y + 1]); |
| 121 | max_v = max_val(max_v, resp_in[(x + 1) * idim0 + y + 1]); |
| 122 | |
| 123 | // Stores corner to {x,y,resp}_out if it's response is maximum compared |
| 124 | // to its 8-neighborhood and greater or equal minimum response |
| 125 | if (v > max_v && v >= (T)min_resp) { |
| 126 | unsigned idx = atomicAdd(count, 1u); |
| 127 | if (idx < max_corners) { |
| 128 | x_out[idx] = (float)x; |
| 129 | y_out[idx] = (float)y; |
| 130 | resp_out[idx] = (float)v; |
| 131 | } |
| 132 | } |
| 133 | } |
| 134 | } |
| 135 | |
| 136 | __global__ void keep_corners(float* x_out, float* y_out, float* resp_out, |
| 137 | const float* x_in, const float* y_in, |