| 248 | // 3x3x3 pixel neighborhood. |
| 249 | template<typename T> |
| 250 | __global__ void detectExtrema(float* x_out, float* y_out, unsigned* layer_out, |
| 251 | unsigned* counter, CParam<T> dog, |
| 252 | const unsigned max_feat, const float threshold) { |
| 253 | const int dim0 = dog.dims[0]; |
| 254 | const int dim1 = dog.dims[1]; |
| 255 | const int imel = dim0 * dim1; |
| 256 | |
| 257 | const int tid_i = threadIdx.x; |
| 258 | const int tid_j = threadIdx.y; |
| 259 | const int bsz_i = blockDim.x; |
| 260 | const int bsz_j = blockDim.y; |
| 261 | const int i = blockIdx.x * bsz_i + tid_i + IMG_BORDER; |
| 262 | const int j = blockIdx.y * bsz_j + tid_j + IMG_BORDER; |
| 263 | |
| 264 | const int x = tid_i + 1; |
| 265 | const int y = tid_j + 1; |
| 266 | |
| 267 | // One pixel border for each side |
| 268 | const int s_i = bsz_i + 2; |
| 269 | const int s_j = bsz_j + 2; |
| 270 | |
| 271 | SharedMemory<float> shared; |
| 272 | float* shrdMem = shared.getPointer(); |
| 273 | float* s_next = shrdMem; |
| 274 | float* s_center = shrdMem + s_i * s_j; |
| 275 | float* s_prev = shrdMem + s_i * s_j * 2; |
| 276 | |
| 277 | for (int l = 1; l < dog.dims[2] - 1; l++) { |
| 278 | const int s_i_half = s_i / 2; |
| 279 | const int s_j_half = s_j / 2; |
| 280 | if (tid_i < s_i_half && tid_j < s_j_half && i < dim0 - IMG_BORDER + 1 && |
| 281 | j < dim1 - IMG_BORDER + 1) { |
| 282 | SNPTR(tid_j, tid_i) = DPTR(l + 1, j - 1, i - 1); |
| 283 | SCPTR(tid_j, tid_i) = DPTR(l, j - 1, i - 1); |
| 284 | SPPTR(tid_j, tid_i) = DPTR(l - 1, j - 1, i - 1); |
| 285 | |
| 286 | SNPTR(tid_j, tid_i + s_i_half) = |
| 287 | DPTR((l + 1), j - 1, i - 1 + s_i_half); |
| 288 | SCPTR(tid_j, tid_i + s_i_half) = DPTR((l), j - 1, i - 1 + s_i_half); |
| 289 | SPPTR(tid_j, tid_i + s_i_half) = |
| 290 | DPTR((l - 1), j - 1, i - 1 + s_i_half); |
| 291 | |
| 292 | SNPTR(tid_j + s_j_half, tid_i) = |
| 293 | DPTR(l + 1, j - 1 + s_j_half, i - 1); |
| 294 | SCPTR(tid_j + s_j_half, tid_i) = DPTR(l, j - 1 + s_j_half, i - 1); |
| 295 | SPPTR(tid_j + s_j_half, tid_i) = |
| 296 | DPTR(l - 1, j - 1 + s_j_half, i - 1); |
| 297 | |
| 298 | SNPTR(tid_j + s_j_half, tid_i + s_i_half) = |
| 299 | DPTR(l + 1, j - 1 + s_j_half, i - 1 + s_i_half); |
| 300 | SCPTR(tid_j + s_j_half, tid_i + s_i_half) = |
| 301 | DPTR(l, j - 1 + s_j_half, i - 1 + s_i_half); |
| 302 | SPPTR(tid_j + s_j_half, tid_i + s_i_half) = |
| 303 | DPTR(l - 1, j - 1 + s_j_half, i - 1 + s_i_half); |
| 304 | } |
| 305 | __syncthreads(); |
| 306 | |
| 307 | float p = SCPTR(y, x); |
nothing calls this directly
no test coverage detected