| 218 | |
| 219 | template<typename T> |
| 220 | __global__ void extract_orb(unsigned* desc_out, const unsigned n_feat, |
| 221 | float* x_in_out, float* y_in_out, |
| 222 | const float* ori_in, float* size_out, |
| 223 | CParam<T> image, const float scl, |
| 224 | const unsigned patch_size, |
| 225 | cudaTextureObject_t luTable) { |
| 226 | unsigned f = blockDim.x * blockIdx.x + threadIdx.x; |
| 227 | |
| 228 | if (f < n_feat) { |
| 229 | unsigned x = (unsigned)round(x_in_out[f]); |
| 230 | unsigned y = (unsigned)round(y_in_out[f]); |
| 231 | float ori = ori_in[f]; |
| 232 | unsigned size = patch_size; |
| 233 | |
| 234 | unsigned r = ceil(patch_size * sqrt(2.f) / 2.f); |
| 235 | if (x < r || y < r || x >= image.dims[1] - r || y >= image.dims[0] - r) |
| 236 | return; |
| 237 | |
| 238 | // Descriptor fixed at 256 bits for now |
| 239 | // Storing descriptor as a vector of 8 x 32-bit unsigned numbers |
| 240 | for (unsigned i = threadIdx.y; i < 16; i += blockDim.y) { |
| 241 | unsigned v = 0; |
| 242 | |
| 243 | // j < 16 for 256 bits descriptor |
| 244 | for (unsigned j = 0; j < 16; j++) { |
| 245 | // Get position from distribution pattern and values of points |
| 246 | // p1 and p2 |
| 247 | int dist_x = lookup(i * 16 * 4 + j * 4, luTable); |
| 248 | int dist_y = lookup(i * 16 * 4 + j * 4 + 1, luTable); |
| 249 | T p1 = get_pixel(x, y, ori, size, dist_x, dist_y, image, |
| 250 | patch_size); |
| 251 | |
| 252 | dist_x = lookup(i * 16 * 4 + j * 4 + 2, luTable); |
| 253 | dist_y = lookup(i * 16 * 4 + j * 4 + 3, luTable); |
| 254 | T p2 = get_pixel(x, y, ori, size, dist_x, dist_y, image, |
| 255 | patch_size); |
| 256 | |
| 257 | // Calculate bit based on p1 and p2 and shifts it to correct |
| 258 | // position |
| 259 | v |= (p1 < p2) << (j + 16 * (i % 2)); |
| 260 | } |
| 261 | |
| 262 | // Store 16 bits of descriptor |
| 263 | atomicAdd(&desc_out[f * 8 + i / 2], v); |
| 264 | } |
| 265 | |
| 266 | if (threadIdx.y == 0) { |
| 267 | x_in_out[f] = round(x * scl); |
| 268 | y_in_out[f] = round(y * scl); |
| 269 | size_out[f] = patch_size * scl; |
| 270 | } |
| 271 | } |
| 272 | } |
| 273 | |
| 274 | template<typename T, typename convAccT> |
| 275 | void orb(unsigned* out_feat, float** d_x, float** d_y, float** d_score, |