src/core/CL/cl_kernels/hog.cl

   1 /*
   2  * Copyright (c) 2017, 2018 ARM Limited.
   3  *
   4  * SPDX-License-Identifier: MIT
   5  *
   6  * Permission is hereby granted, free of charge, to any person obtaining a copy
   7  * of this software and associated documentation files (the "Software"), to
   8  * deal in the Software without restriction, including without limitation the
   9  * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
  10  * sell copies of the Software, and to permit persons to whom the Software is
  11  * furnished to do so, subject to the following conditions:
  12  *
  13  * The above copyright notice and this permission notice shall be included in all
  14  * copies or substantial portions of the Software.
  15  *
  16  * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
  17  * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
  18  * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
  19  * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
  20  * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
  21  * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
  22  * SOFTWARE.
  23  */
  24 #include "helpers.h"
  25 #include "types.h"
  26
  27 #if defined(CELL_WIDTH) && defined(CELL_HEIGHT) && defined(NUM_BINS) && defined(PHASE_SCALE)
  28
  29 /** This OpenCL kernel computes the HOG orientation binning
  30  *
  31  * @attention The following variables must be passed at compile time:
  32  *
  33  * -# -DCELL_WIDTH = Width of the cell
  34  * -# -DCELL_HEIGHT = height of the cell
  35  * -# -DNUM_BINS = Number of bins for each cell
  36  * -# -DPHASE_SCALE = Scale factor used to evaluate the index of the local HOG
  37  *
  38  * @note Each work-item computes a single cell
  39  *
  40  * @param[in]  mag_ptr                             Pointer to the source image which stores the magnitude of the gradient for each pixel. Supported data types: S16
  41  * @param[in]  mag_stride_x                        Stride of the magnitude image in X dimension (in bytes)
  42  * @param[in]  mag_step_x                          mag_stride_x * number of elements along X processed per workitem(in bytes)
  43  * @param[in]  mag_stride_y                        Stride of the magnitude image in Y dimension (in bytes)
  44  * @param[in]  mag_step_y                          mag_stride_y * number of elements along Y processed per workitem(in bytes)
  45  * @param[in]  mag_offset_first_element_in_bytes   The offset of the first element in the magnitude image
  46  * @param[in]  phase_ptr                           Pointer to the source image which stores the phase of the gradient for each pixel. Supported data types: U8
  47  * @param[in]  phase_stride_x                      Stride of the phase image in X dimension (in bytes)
  48  * @param[in]  phase_step_x                        phase_stride_x * number of elements along X processed per workitem(in bytes)
  49  * @param[in]  phase_stride_y                      Stride of the the phase image in Y dimension (in bytes)
  50  * @param[in]  phase_step_y                        phase_stride_y * number of elements along Y processed per workitem(in bytes)
  51  * @param[in]  phase_offset_first_element_in_bytes The offset of the first element in the the phase image
  52  * @param[out] dst_ptr                             Pointer to the destination image which stores the local HOG for each cell Supported data types: F32. Number of channels supported: equal to the number of histogram bins per cell
  53  * @param[in]  dst_stride_x                        Stride of the destination image in X dimension (in bytes)
  54  * @param[in]  dst_step_x                          dst_stride_x * number of elements along X processed per workitem(in bytes)
  55  * @param[in]  dst_stride_y                        Stride of the destination image in Y dimension (in bytes)
  56  * @param[in]  dst_step_y                          dst_stride_y * number of elements along Y processed per workitem(in bytes)
  57  * @param[in]  dst_offset_first_element_in_bytes   The offset of the first element in the destination image
  58  */
  59 __kernel void hog_orientation_binning(IMAGE_DECLARATION(mag),
  60                                       IMAGE_DECLARATION(phase),
  61                                       IMAGE_DECLARATION(dst))
  62 {
  63     float bins[NUM_BINS] = { 0 };
  64
  65     // Compute address for the magnitude and phase images
  66     Image mag   = CONVERT_TO_IMAGE_STRUCT(mag);
  67     Image phase = CONVERT_TO_IMAGE_STRUCT(phase);
  68
  69     __global uchar *mag_row_ptr   = mag.ptr;
  70     __global uchar *phase_row_ptr = phase.ptr;
  71
  72     for(int yc = 0; yc < CELL_HEIGHT; ++yc)
  73     {
  74         int xc = 0;
  75         for(; xc <= (CELL_WIDTH - 4); xc += 4)
  76         {
  77             // Load magnitude and phase values
  78             const float4 mag_f32   = convert_float4(vload4(0, (__global short *)mag_row_ptr + xc));
  79             float4       phase_f32 = convert_float4(vload4(0, phase_row_ptr + xc));
  80
  81             // Scale phase: phase * scale + 0.5f
  82             phase_f32 = (float4)0.5f + phase_f32 * (float4)PHASE_SCALE;
  83
  84             // Compute histogram index.
  85             int4 hidx_s32 = convert_int4(phase_f32);
  86
  87             // Compute magnitude weights (w0 and w1)
  88             const float4 hidx_f32 = convert_float4(hidx_s32);
  89
  90             // w1 = phase_f32 - hidx_s32
  91             const float4 w1_f32 = phase_f32 - hidx_f32;
  92
  93             // w0 = 1.0 - w1
  94             const float4 w0_f32 = (float4)1.0f - w1_f32;
  95
  96             // Calculate the weights for splitting vote
  97             const float4 mag_w0_f32 = mag_f32 * w0_f32;
  98             const float4 mag_w1_f32 = mag_f32 * w1_f32;
  99
 100             // Weighted vote between 2 bins
 101
 102             // Check if the histogram index is equal to NUM_BINS. If so, replace the index with 0
 103             hidx_s32 = select(hidx_s32, (int4)0, hidx_s32 == (int4)(NUM_BINS));
 104
 105             // Bin 0
 106             bins[hidx_s32.s0] += mag_w0_f32.s0;
 107             bins[hidx_s32.s1] += mag_w0_f32.s1;
 108             bins[hidx_s32.s2] += mag_w0_f32.s2;
 109             bins[hidx_s32.s3] += mag_w0_f32.s3;
 110
 111             hidx_s32 += (int4)1;
 112
 113             // Check if the histogram index is equal to NUM_BINS. If so, replace the index with 0
 114             hidx_s32 = select(hidx_s32, (int4)0, hidx_s32 == (int4)(NUM_BINS));
 115
 116             // Bin1
 117             bins[hidx_s32.s0] += mag_w1_f32.s0;
 118             bins[hidx_s32.s1] += mag_w1_f32.s1;
 119             bins[hidx_s32.s2] += mag_w1_f32.s2;
 120             bins[hidx_s32.s3] += mag_w1_f32.s3;
 121         }
 122
 123         // Left over computation
 124         for(; xc < CELL_WIDTH; xc++)
 125         {
 126             const float mag_value   = *((__global short *)mag_row_ptr + xc);
 127             const float phase_value = *(phase_row_ptr + xc) * (float)PHASE_SCALE + 0.5f;
 128             const float w1          = phase_value - floor(phase_value);
 129
 130             // The quantised phase is the histogram index [0, NUM_BINS - 1]
 131             // Check limit of histogram index. If hidx == NUM_BINS, hidx = 0
 132             const uint hidx = (uint)(phase_value) % NUM_BINS;
 133
 134             // Weighted vote between 2 bins
 135             bins[hidx] += mag_value * (1.0f - w1);
 136             bins[(hidx + 1) % NUM_BINS] += mag_value * w1;
 137         }
 138
 139         // Point to the next row of magnitude and phase images
 140         mag_row_ptr += mag_stride_y;
 141         phase_row_ptr += phase_stride_y;
 142     }
 143
 144     // Compute address for the destination image
 145     Image dst = CONVERT_TO_IMAGE_STRUCT(dst);
 146
 147     // Store the local HOG in the global memory
 148     int xc = 0;
 149     for(; xc <= (NUM_BINS - 4); xc += 4)
 150     {
 151         float4 values = vload4(0, bins + xc);
 152
 153         vstore4(values, 0, ((__global float *)dst.ptr) + xc);
 154     }
 155
 156     // Left over stores
 157     for(; xc < NUM_BINS; ++xc)
 158     {
 159         ((__global float *)dst.ptr)[xc] = bins[xc];
 160     }
 161 }
 162 #endif /* CELL_WIDTH and CELL_HEIGHT and NUM_BINS and PHASE_SCALE */
 163
 164 #if defined(NUM_CELLS_PER_BLOCK_HEIGHT) && defined(NUM_BINS_PER_BLOCK_X) && defined(NUM_BINS_PER_BLOCK) && defined(HOG_NORM_TYPE) && defined(L2_HYST_THRESHOLD)
 165
 166 #ifndef L2_NORM
 167 #error The value of enum class HOGNormType::L2_NORM has not be passed to the OpenCL kernel
 168 #endif /* not L2_NORM */
 169
 170 #ifndef L2HYS_NORM
 171 #error The value of enum class HOGNormType::L2HYS_NORM has not be passed to the OpenCL kernel
 172 #endif /* not L2HYS_NORM */
 173
 174 #ifndef L1_NORM
 175 #error The value of enum class HOGNormType::L1_NORM has not be passed to the OpenCL kernel
 176 #endif /* not L1_NORM */
 177
 178 /** This OpenCL kernel computes the HOG block normalization
 179  *
 180  * @attention The following variables must be passed at compile time:
 181  *
 182  * -# -DNUM_CELLS_PER_BLOCK_HEIGHT = Number of cells for each block
 183  * -# -DNUM_BINS_PER_BLOCK_X = Number of bins for each block along the X direction
 184  * -# -DNUM_BINS_PER_BLOCK = Number of bins for each block
 185  * -# -DHOG_NORM_TYPE = Normalization type
 186  * -# -DL2_HYST_THRESHOLD = Threshold used for L2HYS_NORM normalization method
 187  * -# -DL2_NORM = Value of the enum class HOGNormType::L2_NORM
 188  * -# -DL2HYS_NORM = Value of the enum class HOGNormType::L2HYS_NORM
 189  * -# -DL1_NORM = Value of the enum class HOGNormType::L1_NORM
 190  *
 191  * @note Each work-item computes a single block
 192  *
 193  * @param[in]  src_ptr                           Pointer to the source image which stores the local HOG. Supported data types: F32. Number of channels supported: equal to the number of histogram bins per cell
 194  * @param[in]  src_stride_x                      Stride of the source image in X dimension (in bytes)
 195  * @param[in]  src_step_x                        src_stride_x * number of elements along X processed per workitem(in bytes)
 196  * @param[in]  src_stride_y                      Stride of the source image in Y dimension (in bytes)
 197  * @param[in]  src_step_y                        src_stride_y * number of elements along Y processed per workitem(in bytes)
 198  * @param[in]  src_offset_first_element_in_bytes The offset of the first element in the source image
 199  * @param[out] dst_ptr                           Pointer to the destination image which stores the normlized HOG Supported data types: F32. Number of channels supported: equal to the number of histogram bins per block
 200  * @param[in]  dst_stride_x                      Stride of the destination image in X dimension (in bytes)
 201  * @param[in]  dst_step_x                        dst_stride_x * number of elements along X processed per workitem(in bytes)
 202  * @param[in]  dst_stride_y                      Stride of the destination image in Y dimension (in bytes)
 203  * @param[in]  dst_step_y                        dst_stride_y * number of elements along Y processed per workitem(in bytes)
 204  * @param[in]  dst_offset_first_element_in_bytes The offset of the first element in the destination image
 205  */
 206 __kernel void hog_block_normalization(IMAGE_DECLARATION(src),
 207                                       IMAGE_DECLARATION(dst))
 208 {
 209     float  sum     = 0.0f;
 210     float4 sum_f32 = (float4)(0.0f);
 211
 212     // Compute address for the source and destination tensor
 213     Image src = CONVERT_TO_IMAGE_STRUCT(src);
 214     Image dst = CONVERT_TO_IMAGE_STRUCT(dst);
 215
 216     for(size_t yc = 0; yc < NUM_CELLS_PER_BLOCK_HEIGHT; ++yc)
 217     {
 218         const __global float *hist_ptr = (__global float *)(src.ptr + yc * src_stride_y);
 219
 220         int xc = 0;
 221         for(; xc <= (NUM_BINS_PER_BLOCK_X - 16); xc += 16)
 222         {
 223             const float4 val0 = vload4(0, hist_ptr + xc + 0);
 224             const float4 val1 = vload4(0, hist_ptr + xc + 4);
 225             const float4 val2 = vload4(0, hist_ptr + xc + 8);
 226             const float4 val3 = vload4(0, hist_ptr + xc + 12);
 227
 228 #if(HOG_NORM_TYPE == L2_NORM) || (HOG_NORM_TYPE == L2HYS_NORM)
 229             // Compute val^2 for L2_NORM or L2HYS_NORM
 230             sum_f32 += val0 * val0;
 231             sum_f32 += val1 * val1;
 232             sum_f32 += val2 * val2;
 233             sum_f32 += val3 * val3;
 234 #else  /* (HOG_NORM_TYPE == L2_NORM) || (HOG_NORM_TYPE == L2HYS_NORM) */
 235             // Compute |val| for L1_NORM
 236             sum_f32 += fabs(val0);
 237             sum_f32 += fabs(val1);
 238             sum_f32 += fabs(val2);
 239             sum_f32 += fabs(val3);
 240 #endif /* (HOG_NORM_TYPE == L2_NORM) || (HOG_NORM_TYPE == L2HYS_NORM) */
 241
 242             // Store linearly the input values un-normalized in the output image. These values will be reused for the normalization.
 243             // This approach will help us to be cache friendly in the next for loop where the normalization will be done because all the values
 244             // will be accessed consecutively
 245             vstore4(val0, 0, ((__global float *)dst.ptr) + xc + 0 + yc * NUM_BINS_PER_BLOCK_X);
 246             vstore4(val1, 0, ((__global float *)dst.ptr) + xc + 4 + yc * NUM_BINS_PER_BLOCK_X);
 247             vstore4(val2, 0, ((__global float *)dst.ptr) + xc + 8 + yc * NUM_BINS_PER_BLOCK_X);
 248             vstore4(val3, 0, ((__global float *)dst.ptr) + xc + 12 + yc * NUM_BINS_PER_BLOCK_X);
 249         }
 250
 251         // Compute left over
 252         for(; xc < NUM_BINS_PER_BLOCK_X; ++xc)
 253         {
 254             const float val = hist_ptr[xc];
 255
 256 #if(HOG_NORM_TYPE == L2_NORM) || (HOG_NORM_TYPE == L2HYS_NORM)
 257             sum += val * val;
 258 #else  /* (HOG_NORM_TYPE == L2_NORM) || (HOG_NORM_TYPE == L2HYS_NORM) */
 259             sum += fabs(val);
 260 #endif /* (HOG_NORM_TYPE == L2_NORM) || (HOG_NORM_TYPE == L2HYS_NORM) */
 261
 262             ((__global float *)dst.ptr)[xc + 0 + yc * NUM_BINS_PER_BLOCK_X] = val;
 263         }
 264     }
 265
 266     sum += dot(sum_f32, (float4)1.0f);
 267
 268     float scale = 1.0f / (sqrt(sum) + NUM_BINS_PER_BLOCK * 0.1f);
 269
 270 #if(HOG_NORM_TYPE == L2HYS_NORM)
 271     // Reset sum
 272     sum_f32 = (float4)0.0f;
 273     sum     = 0.0f;
 274
 275     int k = 0;
 276     for(; k <= NUM_BINS_PER_BLOCK - 16; k += 16)
 277     {
 278         float4 val0 = vload4(0, ((__global float *)dst.ptr) + k + 0);
 279         float4 val1 = vload4(0, ((__global float *)dst.ptr) + k + 4);
 280         float4 val2 = vload4(0, ((__global float *)dst.ptr) + k + 8);
 281         float4 val3 = vload4(0, ((__global float *)dst.ptr) + k + 12);
 282
 283         // Scale val
 284         val0 = val0 * (float4)scale;
 285         val1 = val1 * (float4)scale;
 286         val2 = val2 * (float4)scale;
 287         val3 = val3 * (float4)scale;
 288
 289         // Clip val if over _threshold_l2hys
 290         val0 = fmin(val0, (float4)L2_HYST_THRESHOLD);
 291         val1 = fmin(val1, (float4)L2_HYST_THRESHOLD);
 292         val2 = fmin(val2, (float4)L2_HYST_THRESHOLD);
 293         val3 = fmin(val3, (float4)L2_HYST_THRESHOLD);
 294
 295         // Compute val^2
 296         sum_f32 += val0 * val0;
 297         sum_f32 += val1 * val1;
 298         sum_f32 += val2 * val2;
 299         sum_f32 += val3 * val3;
 300
 301         vstore4(val0, 0, ((__global float *)dst.ptr) + k + 0);
 302         vstore4(val1, 0, ((__global float *)dst.ptr) + k + 4);
 303         vstore4(val2, 0, ((__global float *)dst.ptr) + k + 8);
 304         vstore4(val3, 0, ((__global float *)dst.ptr) + k + 12);
 305     }
 306
 307     // Compute left over
 308     for(; k < NUM_BINS_PER_BLOCK; ++k)
 309     {
 310         float val = ((__global float *)dst.ptr)[k] * scale;
 311
 312         // Clip scaled input_value if over L2_HYST_THRESHOLD
 313         val = fmin(val, (float)L2_HYST_THRESHOLD);
 314
 315         sum += val * val;
 316
 317         ((__global float *)dst.ptr)[k] = val;
 318     }
 319
 320     sum += dot(sum_f32, (float4)1.0f);
 321
 322     // We use the same constants of OpenCV
 323     scale = 1.0f / (sqrt(sum) + 1e-3f);
 324
 325 #endif /* (HOG_NORM_TYPE == L2HYS_NORM) */
 326
 327     int i = 0;
 328     for(; i <= (NUM_BINS_PER_BLOCK - 16); i += 16)
 329     {
 330         float4 val0 = vload4(0, ((__global float *)dst.ptr) + i + 0);
 331         float4 val1 = vload4(0, ((__global float *)dst.ptr) + i + 4);
 332         float4 val2 = vload4(0, ((__global float *)dst.ptr) + i + 8);
 333         float4 val3 = vload4(0, ((__global float *)dst.ptr) + i + 12);
 334
 335         // Multiply val by the normalization scale factor
 336         val0 = val0 * (float4)scale;
 337         val1 = val1 * (float4)scale;
 338         val2 = val2 * (float4)scale;
 339         val3 = val3 * (float4)scale;
 340
 341         vstore4(val0, 0, ((__global float *)dst.ptr) + i + 0);
 342         vstore4(val1, 0, ((__global float *)dst.ptr) + i + 4);
 343         vstore4(val2, 0, ((__global float *)dst.ptr) + i + 8);
 344         vstore4(val3, 0, ((__global float *)dst.ptr) + i + 12);
 345     }
 346
 347     for(; i < NUM_BINS_PER_BLOCK; ++i)
 348     {
 349         ((__global float *)dst.ptr)[i] *= scale;
 350     }
 351 }
 352 #endif /* NUM_CELLS_PER_BLOCK_HEIGHT and NUM_BINS_PER_BLOCK_X and NUM_BINS_PER_BLOCK and HOG_NORM_TYPE and L2_HYST_THRESHOLD */
 353
 354 #if defined(NUM_BLOCKS_PER_DESCRIPTOR_Y) && defined(NUM_BINS_PER_DESCRIPTOR_X) && defined(THRESHOLD) && defined(MAX_NUM_DETECTION_WINDOWS) && defined(IDX_CLASS) && defined(BLOCK_STRIDE_WIDTH) && defined(BLOCK_STRIDE_HEIGHT) && defined(DETECTION_WINDOW_WIDTH) && defined(DETECTION_WINDOW_HEIGHT)
 355
 356 /** This OpenCL kernel computes the HOG detector using linear SVM
 357  *
 358  * @attention The following variables must be passed at compile time:
 359  *
 360  * -# -DNUM_BLOCKS_PER_DESCRIPTOR_Y = Number of blocks per descriptor along the Y direction
 361  * -# -DNUM_BINS_PER_DESCRIPTOR_X = Number of bins per descriptor along the X direction
 362  * -# -DTHRESHOLD = Threshold for the distance between features and SVM classifying plane
 363  * -# -DMAX_NUM_DETECTION_WINDOWS = Maximum number of possible detection windows. It is equal to the size of the DetectioWindow array
 364  * -# -DIDX_CLASS = Index of the class to detect
 365  * -# -DBLOCK_STRIDE_WIDTH = Block stride for the X direction
 366  * -# -DBLOCK_STRIDE_HEIGHT = Block stride for the Y direction
 367  * -# -DDETECTION_WINDOW_WIDTH = Width of the detection window
 368  * -# -DDETECTION_WINDOW_HEIGHT = Height of the detection window
 369  *
 370  * @note Each work-item computes a single detection window
 371  *
 372  * @param[in]  src_ptr                           Pointer to the source image which stores the local HOG. Supported data types: F32. Number of channels supported: equal to the number of histogram bins per cell
 373  * @param[in]  src_stride_x                      Stride of the source image in X dimension (in bytes)
 374  * @param[in]  src_step_x                        src_stride_x * number of elements along X processed per workitem(in bytes)
 375  * @param[in]  src_stride_y                      Stride of the source image in Y dimension (in bytes)
 376  * @param[in]  src_step_y                        src_stride_y * number of elements along Y processed per workitem(in bytes)
 377  * @param[in]  src_offset_first_element_in_bytes The offset of the first element in the source image
 378  * @param[in]  hog_descriptor                    Pointer to HOG descriptor. Supported data types: F32
 379  * @param[out] dst                               Pointer to DetectionWindow array
 380  * @param[out] num_detection_windows             Number of objects detected
 381  */
 382 __kernel void hog_detector(IMAGE_DECLARATION(src),
 383                            __global float *hog_descriptor,
 384                            __global DetectionWindow *dst,
 385                            __global uint *num_detection_windows)
 386 {
 387     // Check if the DetectionWindow array is full
 388     if(*num_detection_windows >= MAX_NUM_DETECTION_WINDOWS)
 389     {
 390         return;
 391     }
 392
 393     Image src = CONVERT_TO_IMAGE_STRUCT(src);
 394
 395     const int src_step_y_f32 = src_stride_y / sizeof(float);
 396
 397     // Init score_f32 with 0
 398     float4 score_f32 = (float4)0.0f;
 399
 400     // Init score with 0
 401     float score = 0.0f;
 402
 403     __global float *src_row_ptr = (__global float *)src.ptr;
 404
 405     // Compute Linear SVM
 406     for(int yb = 0; yb < NUM_BLOCKS_PER_DESCRIPTOR_Y; ++yb, src_row_ptr += src_step_y_f32)
 407     {
 408         int xb = 0;
 409
 410         const int offset_y = yb * NUM_BINS_PER_DESCRIPTOR_X;
 411
 412         for(; xb < (int)NUM_BINS_PER_DESCRIPTOR_X - 8; xb += 8)
 413         {
 414             // Load descriptor values
 415             float4 a0_f32 = vload4(0, src_row_ptr + xb + 0);
 416             float4 a1_f32 = vload4(0, src_row_ptr + xb + 4);
 417
 418             float4 b0_f32 = vload4(0, hog_descriptor + xb + 0 + offset_y);
 419             float4 b1_f32 = vload4(0, hog_descriptor + xb + 4 + offset_y);
 420
 421             // Multiply accumulate
 422             score_f32 += a0_f32 * b0_f32;
 423             score_f32 += a1_f32 * b1_f32;
 424         }
 425
 426         for(; xb < NUM_BINS_PER_DESCRIPTOR_X; ++xb)
 427         {
 428             const float a = src_row_ptr[xb];
 429             const float b = hog_descriptor[xb + offset_y];
 430
 431             score += a * b;
 432         }
 433     }
 434
 435     score += dot(score_f32, (float4)1.0f);
 436
 437     // Add the bias. The bias is located at the position (descriptor_size() - 1)
 438     // (descriptor_size - 1) = NUM_BINS_PER_DESCRIPTOR_X * NUM_BLOCKS_PER_DESCRIPTOR_Y
 439     score += hog_descriptor[NUM_BINS_PER_DESCRIPTOR_X * NUM_BLOCKS_PER_DESCRIPTOR_Y];
 440
 441     if(score > (float)THRESHOLD)
 442     {
 443         int id = atomic_inc(num_detection_windows);
 444         if(id < MAX_NUM_DETECTION_WINDOWS)
 445         {
 446             dst[id].x         = get_global_id(0) * BLOCK_STRIDE_WIDTH;
 447             dst[id].y         = get_global_id(1) * BLOCK_STRIDE_HEIGHT;
 448             dst[id].width     = DETECTION_WINDOW_WIDTH;
 449             dst[id].height    = DETECTION_WINDOW_HEIGHT;
 450             dst[id].idx_class = IDX_CLASS;
 451             dst[id].score     = score;
 452         }
 453     }
 454 }
 455 #endif /* NUM_BLOCKS_PER_DESCRIPTOR_Y && NUM_BINS_PER_DESCRIPTOR_X && THRESHOLD && MAX_NUM_DETECTION_WINDOWS && IDX_CLASS &&
 456         * BLOCK_STRIDE_WIDTH && BLOCK_STRIDE_HEIGHT && DETECTION_WINDOW_WIDTH && DETECTION_WINDOW_HEIGHT */