| 190 | } |
| 191 | |
| 192 | __inline__ __device__ void normalizeGLOHDesc(float* desc, float* accum, |
| 193 | const int histlen) { |
| 194 | int tid_x = threadIdx.x; |
| 195 | int tid_y = threadIdx.y; |
| 196 | int bsz_x = blockDim.x; |
| 197 | |
| 198 | for (int i = tid_x; i < histlen; i += bsz_x) |
| 199 | accum[i] = desc[tid_y * histlen + i] * desc[tid_y * histlen + i]; |
| 200 | __syncthreads(); |
| 201 | |
| 202 | if (tid_x < 128) accum[tid_x] += accum[tid_x + 128]; |
| 203 | __syncthreads(); |
| 204 | if (tid_x < 64) accum[tid_x] += accum[tid_x + 64]; |
| 205 | __syncthreads(); |
| 206 | if (tid_x < 32) accum[tid_x] += accum[tid_x + 32]; |
| 207 | __syncthreads(); |
| 208 | if (tid_x < 16) |
| 209 | // GLOH is 272-dimensional, accumulating last 16 descriptors |
| 210 | accum[tid_x] += accum[tid_x + 16] + accum[tid_x + 256]; |
| 211 | __syncthreads(); |
| 212 | if (tid_x < 8) accum[tid_x] += accum[tid_x + 8]; |
| 213 | __syncthreads(); |
| 214 | if (tid_x < 4) accum[tid_x] += accum[tid_x + 4]; |
| 215 | __syncthreads(); |
| 216 | if (tid_x < 2) accum[tid_x] += accum[tid_x + 2]; |
| 217 | __syncthreads(); |
| 218 | if (tid_x < 1) accum[tid_x] += accum[tid_x + 1]; |
| 219 | __syncthreads(); |
| 220 | |
| 221 | float len_sq = accum[0]; |
| 222 | float len_inv = 1.0f / sqrtf(len_sq); |
| 223 | |
| 224 | for (int i = tid_x; i < histlen; i += bsz_x) { |
| 225 | desc[tid_y * histlen + i] *= len_inv; |
| 226 | } |
| 227 | __syncthreads(); |
| 228 | } |
| 229 | |
| 230 | template<typename T> |
| 231 | __global__ void sub(Param<T> out, CParam<T> in, const unsigned nel, |
no outgoing calls
no test coverage detected