src/core/CL/cl_kernels/fast_corners.cl

   1 /*
   2  * Copyright (c) 2016-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 /* The map table to retrieve the 16 texels in the Bresenham circle of radius 3 with center in P.
  28  *
  29  *      . . F 0 1 . . .
  30  *      . E . . . 2 . .
  31  *      D . . . . . 3 .
  32  *      C . . P . . 4 .
  33  *      B . . . . . 5 .
  34  *      . A . . . 6 . .
  35  *      . . 9 8 7 . . .
  36  */
  37 constant int offsets_s[16][2] =
  38 {
  39     { 0, -3 },  // 0
  40     { 1, -3 },  // 1
  41     { 2, -2 },  // 2
  42     { 3, -1 },  // 3
  43     { 3, 0 },   // 4
  44     { 3, 1 },   // 5
  45     { 2, 2 },   // 6
  46     { 1, 3 },   // 7
  47     { 0, 3 },   // 8
  48     { -1, 3 },  // 9
  49     { -2, 2 },  // A
  50     { -3, 1 },  // B
  51     { -3, 0 },  // C
  52     { -3, -1 }, // D
  53     { -2, -2 }, // E
  54     { -1, -3 }, // F
  55 };
  56
  57 /** Load a pixel and set the mask values.
  58  *
  59  * @param[in]  ptr         The pointer to the starting address of source image
  60  * @param[in]  a           Index to indicate the position in the Bresenham circle
  61  * @param[in]  stride      Stride of source image in x dimension
  62  * @param[in]  dark        The left end of the threshold range
  63  * @param[in]  bright      The right end of the threshold range
  64  * @param[out] dark_mask   The bit-set mask records dark pixels. Its bit is set as 1 if the corresponding pixel is dark
  65  * @param[out] bright_mask The bit-set mask records bright pixels. Its bit is set as 1 if the corresponding pixel is bright
  66  *
  67  */
  68 #define LOAD_AND_SET_MASK(ptr, a, stride, dark, bright, dark_mask, bright_mask) \
  69     {                                                                           \
  70         unsigned char pixel;                                                    \
  71         pixel = *(ptr + (int)stride * offsets_s[a][1] + offsets_s[a][0]);       \
  72         dark_mask |= (pixel < dark) << a;                                       \
  73         bright_mask |= (pixel > bright) << a;                                   \
  74     }
  75
  76 /** Checks if a pixel is a corner. Pixel is considerred as a corner if the 9 continuous pixels in the Bresenham circle are bright or dark.
  77  *
  78  * @param[in]  bright_mask The mask recording postions of bright pixels
  79  * @param[in]  dark_mask   The mask recording postions of dark pixels
  80  * @param[out] isCorner    Indicate whether candidate pixel is corner
  81  */
  82 #define CHECK_CORNER(bright_mask, dark_mask, isCorner)    \
  83     {                                                     \
  84         for(int i = 0; i < 16; i++)                       \
  85         {                                                 \
  86             isCorner |= ((bright_mask & 0x1FF) == 0x1FF); \
  87             isCorner |= ((dark_mask & 0x1FF) == 0x1FF);   \
  88             if(isCorner)                                  \
  89             {                                             \
  90                 break;                                    \
  91             }                                             \
  92             bright_mask >>= 1;                            \
  93             dark_mask >>= 1;                              \
  94         }                                                 \
  95     }
  96
  97 /* Calculate pixel's strength */
  98 uchar compute_strength(uchar candidate_pixel, __global unsigned char *ptr, unsigned int stride, unsigned char threshold)
  99 {
 100     short a = threshold;
 101     short b = 255;
 102     while(b - a > 1)
 103     {
 104         uchar        c           = convert_uchar_sat((a + b) / 2);
 105         unsigned int bright_mask = 0;
 106         unsigned int dark_mask   = 0;
 107
 108         unsigned char p_bright = add_sat(candidate_pixel, c);
 109         unsigned char p_dark   = sub_sat(candidate_pixel, c);
 110
 111         bool isCorner = 0;
 112
 113         for(uint i = 0; i < 16; i++)
 114         {
 115             LOAD_AND_SET_MASK(ptr, i, stride, p_dark, p_bright, dark_mask, bright_mask)
 116         }
 117
 118         bright_mask |= (bright_mask << 16);
 119         dark_mask |= (dark_mask << 16);
 120         CHECK_CORNER(bright_mask, dark_mask, isCorner);
 121
 122         if(isCorner)
 123         {
 124             a = convert_short(c);
 125         }
 126         else
 127         {
 128             b = convert_short(c);
 129         }
 130     }
 131     return a;
 132 }
 133
 134 /** Fast corners implementation. Calculates and returns the strength of each pixel.
 135  *
 136  * The algorithm loops through the 16 pixels in the Bresenham circle and set low 16 bit of masks if corresponding pixel is bright
 137  * or dark. It then copy the low 16 bit to the high 16 bit of the masks. Right shift the bit to check whether the 9 continuous bits
 138  * from the LSB are set.
 139  *
 140  * @param[in]  input_ptr                            Pointer to the first source image. Supported data types: U8
 141  * @param[in]  input_stride_x                       Stride of the first source image in X dimension (in bytes)
 142  * @param[in]  input_step_x                         input_stride_x * number of elements along X processed per workitem(in bytes)
 143  * @param[in]  input_stride_y                       Stride of the first source image in Y dimension (in bytes)
 144  * @param[in]  input_step_y                         input_stride_y * number of elements along Y processed per workitem(in bytes)
 145  * @param[in]  input_offset_first_element_in_bytes  The offset of the first element in the first source image
 146  * @param[out] output_ptr                           Pointer to the first source image. Supported data types: U8
 147  * @param[in]  output_stride_x                      Stride of the first source image in X dimension (in bytes)
 148  * @param[in]  output_step_x                        output_stride_x * number of elements along X processed per workitem(in bytes)
 149  * @param[in]  output_stride_y                      Stride of the first source image in Y dimension (in bytes)
 150  * @param[in]  output_step_y                        output_stride_y * number of elements along Y processed per workitem(in bytes)
 151  * @param[in]  output_offset_first_element_in_bytes The offset of the first element in the first source image
 152  * @param[in]  threshold_value                      Threshold value.
 153  *
 154  */
 155 __kernel void fast_corners(
 156     IMAGE_DECLARATION(input),
 157     IMAGE_DECLARATION(output),
 158     float threshold_value)
 159 {
 160     Image in  = CONVERT_TO_IMAGE_STRUCT(input);
 161     Image out = CONVERT_TO_IMAGE_STRUCT(output);
 162
 163     const unsigned char threshold = (uchar)threshold_value;
 164
 165     unsigned int bright_mask = 0;
 166     unsigned int dark_mask   = 0;
 167
 168     unsigned char isCorner = 0;
 169
 170     unsigned char p        = *in.ptr;
 171     unsigned char p_bright = add_sat(p, threshold);
 172     unsigned char p_dark   = sub_sat(p, threshold);
 173
 174     LOAD_AND_SET_MASK(in.ptr, 0, input_stride_y, p_dark, p_bright, dark_mask, bright_mask)
 175     LOAD_AND_SET_MASK(in.ptr, 4, input_stride_y, p_dark, p_bright, dark_mask, bright_mask)
 176     LOAD_AND_SET_MASK(in.ptr, 8, input_stride_y, p_dark, p_bright, dark_mask, bright_mask)
 177     LOAD_AND_SET_MASK(in.ptr, 12, input_stride_y, p_dark, p_bright, dark_mask, bright_mask)
 178
 179     if(((bright_mask | dark_mask) & 0x1111) == 0)
 180     {
 181         *out.ptr = 0;
 182         return;
 183     }
 184
 185     LOAD_AND_SET_MASK(in.ptr, 1, input_stride_y, p_dark, p_bright, dark_mask, bright_mask)
 186     LOAD_AND_SET_MASK(in.ptr, 2, input_stride_y, p_dark, p_bright, dark_mask, bright_mask)
 187     LOAD_AND_SET_MASK(in.ptr, 3, input_stride_y, p_dark, p_bright, dark_mask, bright_mask)
 188     LOAD_AND_SET_MASK(in.ptr, 5, input_stride_y, p_dark, p_bright, dark_mask, bright_mask)
 189     LOAD_AND_SET_MASK(in.ptr, 6, input_stride_y, p_dark, p_bright, dark_mask, bright_mask)
 190     LOAD_AND_SET_MASK(in.ptr, 7, input_stride_y, p_dark, p_bright, dark_mask, bright_mask)
 191     LOAD_AND_SET_MASK(in.ptr, 9, input_stride_y, p_dark, p_bright, dark_mask, bright_mask)
 192     LOAD_AND_SET_MASK(in.ptr, 10, input_stride_y, p_dark, p_bright, dark_mask, bright_mask)
 193     LOAD_AND_SET_MASK(in.ptr, 11, input_stride_y, p_dark, p_bright, dark_mask, bright_mask)
 194     LOAD_AND_SET_MASK(in.ptr, 13, input_stride_y, p_dark, p_bright, dark_mask, bright_mask)
 195     LOAD_AND_SET_MASK(in.ptr, 14, input_stride_y, p_dark, p_bright, dark_mask, bright_mask)
 196     LOAD_AND_SET_MASK(in.ptr, 15, input_stride_y, p_dark, p_bright, dark_mask, bright_mask)
 197
 198     bright_mask |= (bright_mask << 16);
 199     dark_mask |= (dark_mask << 16);
 200
 201     CHECK_CORNER(bright_mask, dark_mask, isCorner)
 202
 203     if(!isCorner)
 204     {
 205         *out.ptr = 0;
 206         return;
 207     }
 208
 209 #ifdef USE_MAXSUPPRESSION
 210     *out.ptr = compute_strength(p, in.ptr, input_stride_y, threshold);
 211 #else  /* USE_MAXSUPPRESSION */
 212     *out.ptr = 1;
 213 #endif /* USE_MAXSUPPRESSION */
 214 }
 215
 216 /** Copy result to Keypoint buffer and count number of corners
 217  *
 218  * @param[in]  input_ptr                           Pointer to the image with calculated strenghs. Supported data types: U8
 219  * @param[in]  input_stride_x                      Stride of the first source image in X dimension (in bytes)
 220  * @param[in]  input_step_x                        input_stride_x * number of elements along X processed per workitem(in bytes)
 221  * @param[in]  input_stride_y                      Stride of the first source image in Y dimension (in bytes)
 222  * @param[in]  input_step_y                        input_stride_y * number of elements along Y processed per workitem(in bytes)
 223  * @param[in]  input_offset_first_element_in_bytes The offset of the first element in the first source image
 224  * @param[in]  max_num_points                      The maximum number of keypoints the array can hold
 225  * @param[out] offset                              The number of skipped pixels in x dimension
 226  * @param[out] num_of_points                       Number of points found
 227  * @param[out] out                                 The keypoints found
 228  *
 229  */
 230 __kernel void copy_to_keypoint(
 231     IMAGE_DECLARATION(input),
 232     uint     max_num_points,
 233     uint     offset,
 234     __global uint *num_of_points,
 235     __global Keypoint *out)
 236 {
 237 #ifndef UPDATE_NUMBER
 238     if(*num_of_points >= max_num_points)
 239     {
 240         return;
 241     }
 242 #endif /* UPDATE_NUMBER */
 243
 244     Image in = CONVERT_TO_IMAGE_STRUCT(input);
 245
 246     uchar value = *in.ptr;
 247
 248     if(value > 0)
 249     {
 250         int id = atomic_inc(num_of_points);
 251         if(id < max_num_points)
 252         {
 253             out[id].strength        = value;
 254             out[id].x               = get_global_id(0) + offset;
 255             out[id].y               = get_global_id(1) + offset;
 256             out[id].tracking_status = 1;
 257             out[id].scale           = 0.f;
 258             out[id].orientation     = 0.f;
 259             out[id].error           = 0.f;
 260         }
 261     }
 262 }