| 156 | } |
| 157 | |
| 158 | __inline__ __device__ void normalizeDesc(float* desc, float* accum, |
| 159 | const int histlen) { |
| 160 | int tid_x = threadIdx.x; |
| 161 | int tid_y = threadIdx.y; |
| 162 | int bsz_x = blockDim.x; |
| 163 | |
| 164 | for (int i = tid_x; i < histlen; i += bsz_x) |
| 165 | accum[i] = desc[tid_y * histlen + i] * desc[tid_y * histlen + i]; |
| 166 | __syncthreads(); |
| 167 | |
| 168 | if (tid_x < 64) accum[tid_x] += accum[tid_x + 64]; |
| 169 | __syncthreads(); |
| 170 | if (tid_x < 32) accum[tid_x] += accum[tid_x + 32]; |
| 171 | __syncthreads(); |
| 172 | if (tid_x < 16) accum[tid_x] += accum[tid_x + 16]; |
| 173 | __syncthreads(); |
| 174 | if (tid_x < 8) accum[tid_x] += accum[tid_x + 8]; |
| 175 | __syncthreads(); |
| 176 | if (tid_x < 4) accum[tid_x] += accum[tid_x + 4]; |
| 177 | __syncthreads(); |
| 178 | if (tid_x < 2) accum[tid_x] += accum[tid_x + 2]; |
| 179 | __syncthreads(); |
| 180 | if (tid_x < 1) accum[tid_x] += accum[tid_x + 1]; |
| 181 | __syncthreads(); |
| 182 | |
| 183 | float len_sq = accum[0]; |
| 184 | float len_inv = 1.0f / sqrtf(len_sq); |
| 185 | |
| 186 | for (int i = tid_x; i < histlen; i += bsz_x) { |
| 187 | desc[tid_y * histlen + i] *= len_inv; |
| 188 | } |
| 189 | __syncthreads(); |
| 190 | } |
| 191 | |
| 192 | __inline__ __device__ void normalizeGLOHDesc(float* desc, float* accum, |
| 193 | const int histlen) { |