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(DETECTION_WINDOW_STRIDE_WIDTH) && defined(DETECTION_WINDOW_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 * -# -DDETECTION_WINDOW_STRIDE_WIDTH = Detection window stride for the X direction 366 * -# -DDETECTION_WINDOW_STRIDE_HEIGHT = Detection window 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) * DETECTION_WINDOW_STRIDE_WIDTH; 447 dst[id].y = get_global_id(1) * DETECTION_WINDOW_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 * DETECTION_WINDOW_STRIDE_WIDTH && DETECTION_WINDOW_STRIDE_HEIGHT && DETECTION_WINDOW_WIDTH && DETECTION_WINDOW_HEIGHT */ 457