| 192 | |
| 193 | template<typename T, int arc_length> |
| 194 | __global__ void locate_features(CParam<T> in, float *score, const float thr, |
| 195 | const unsigned edge, |
| 196 | cudaTextureObject_t luTable) { |
| 197 | unsigned ix = threadIdx.x; |
| 198 | unsigned iy = threadIdx.y; |
| 199 | unsigned bx = blockDim.x; |
| 200 | unsigned by = blockDim.y; |
| 201 | unsigned x = bx * blockIdx.x + ix + edge; |
| 202 | unsigned y = by * blockIdx.y + iy + edge; |
| 203 | unsigned lx = bx / 2 + 3; |
| 204 | unsigned ly = by / 2 + 3; |
| 205 | |
| 206 | SharedMemory<T> shared; |
| 207 | T *local_image_curr = shared.getPointer(); |
| 208 | load_shared_image(in, local_image_curr, ix, iy, bx, by, x, y, lx, ly, edge); |
| 209 | __syncthreads(); |
| 210 | locate_features_core<T, arc_length>(local_image_curr, score, in.dims[0], |
| 211 | in.dims[1], thr, x, y, edge, luTable); |
| 212 | } |
| 213 | |
| 214 | template<bool nonmax> |
| 215 | __global__ void non_max_counts(unsigned *d_counts, unsigned *d_offsets, |
nothing calls this directly
no test coverage detected