COMPMID-344 Updated doxygen

Change-Id: I32f7b84daa560e460b77216add529c8fa8b327ae

1 files changed, 260 insertions, 0 deletions

diff --git a/src/core/CL/cl_kernels/fast_corners.cl b/src/core/CL/cl_kernels/fast_corners.cl
new file mode 100644
index 0000000000..470d14a7b0
--- /dev/null
+++ b/src/core/CL/cl_kernels/fast_corners.cl
@@ -0,0 +1,260 @@
+/*
+ * Copyright (c) 2016, 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "helpers.h"
+#include "types.h"
+
+/* The map table to retrieve the 16 texels in the Bresenham circle of radius 3 with center in P.
+ *
+ *      . . F 0 1 . . .
+ *      . E . . . 2 . .
+ *      D . . . . . 3 .
+ *      C . . P . . 4 .
+ *      B . . . . . 5 .
+ *      . A . . . 6 . .
+ *      . . 9 8 7 . . .
+ */
+constant int offsets_s[16][2] =
+{
+    { 0, -3 },  // 0
+    { 1, -3 },  // 1
+    { 2, -2 },  // 2
+    { 3, -1 },  // 3
+    { 3, 0 },   // 4
+    { 3, 1 },   // 5
+    { 2, 2 },   // 6
+    { 1, 3 },   // 7
+    { 0, 3 },   // 8
+    { -1, 3 },  // 9
+    { -2, 2 },  // A
+    { -3, 1 },  // B
+    { -3, 0 },  // C
+    { -3, -1 }, // D
+    { -2, -2 }, // E
+    { -1, -3 }, // F
+};
+
+/** Load a pixel and set the mask values.
+ *
+ * @param[in]  ptr         The pointer to the starting address of source image
+ * @param[in]  a           Index to indicate the position in the Bresenham circle
+ * @param[in]  stride      Stride of source image in x dimension
+ * @param[in]  dark        The left end of the threshold range
+ * @param[in]  bright      The right end of the threshold range
+ * @param[out] dark_mask   The bit-set mask records dark pixels. Its bit is set as 1 if the corresponding pixel is dark
+ * @param[out] bright_mask The bit-set mask records bright pixels. Its bit is set as 1 if the corresponding pixel is bright
+ *
+ */
+#define LOAD_AND_SET_MASK(ptr, a, stride, dark, bright, dark_mask, bright_mask) \
+    {                                                                           \
+        unsigned char pixel;                                                    \
+        pixel = *(ptr + (int)stride * offsets_s[a][1] + offsets_s[a][0]);       \
+        dark_mask |= (pixel < dark) << a;                                       \
+        bright_mask |= (pixel > bright) << a;                                   \
+    }
+
+/** 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.
+ *
+ * @param[in]  bright_mask The mask recording postions of bright pixels
+ * @param[in]  dark_mask   The mask recording postions of dark pixels
+ * @param[out] isCorner    Indicate whether candidate pixel is corner
+ */
+#define CHECK_CORNER(bright_mask, dark_mask, isCorner)    \
+    {                                                     \
+        for(int i = 0; i < 16; i++)                       \
+        {                                                 \
+            isCorner |= ((bright_mask & 0x1FF) == 0x1FF); \
+            isCorner |= ((dark_mask & 0x1FF) == 0x1FF);   \
+            if(isCorner)                                  \
+            {                                             \
+                break;                                    \
+            }                                             \
+            bright_mask >>= 1;                            \
+            dark_mask >>= 1;                              \
+        }                                                 \
+    }
+
+/* Calculate pixel's strength */
+uchar compute_strength(uchar candidate_pixel, __global unsigned char *ptr, unsigned int stride, unsigned char threshold)
+{
+    short a = threshold;
+    short b = 255;
+    while(b - a > 1)
+    {
+        uchar        c           = convert_uchar_sat((a + b) / 2);
+        unsigned int bright_mask = 0;
+        unsigned int dark_mask   = 0;
+
+        unsigned char p_bright = add_sat(candidate_pixel, c);
+        unsigned char p_dark   = sub_sat(candidate_pixel, c);
+
+        bool isCorner = 0;
+
+        for(uint i = 0; i < 16; i++)
+        {
+            LOAD_AND_SET_MASK(ptr, i, stride, p_dark, p_bright, dark_mask, bright_mask)
+        }
+
+        bright_mask |= (bright_mask << 16);
+        dark_mask |= (dark_mask << 16);
+        CHECK_CORNER(bright_mask, dark_mask, isCorner);
+
+        if(isCorner)
+        {
+            a = convert_short(c);
+        }
+        else
+        {
+            b = convert_short(c);
+        }
+    }
+    return a;
+}
+
+/** Fast corners implementation. Calculates and returns the strength of each pixel.
+ *
+ * The algorithm loops through the 16 pixels in the Bresenham circle and set low 16 bit of masks if corresponding pixel is bright
+ * 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
+ * from the LSB are set.
+ *
+ * @param[in]  input_ptr                            Pointer to the first source image. Supported data types: U8
+ * @param[in]  input_stride_x                       Stride of the first source image in X dimension (in bytes)
+ * @param[in]  input_step_x                         input_stride_x * number of elements along X processed per workitem(in bytes)
+ * @param[in]  input_stride_y                       Stride of the first source image in Y dimension (in bytes)
+ * @param[in]  input_step_y                         input_stride_y * number of elements along Y processed per workitem(in bytes)
+ * @param[in]  input_offset_first_element_in_bytes  The offset of the first element in the first source image
+ * @param[out] output_ptr                           Pointer to the first source image. Supported data types: U8
+ * @param[in]  output_stride_x                      Stride of the first source image in X dimension (in bytes)
+ * @param[in]  output_step_x                        output_stride_x * number of elements along X processed per workitem(in bytes)
+ * @param[in]  output_stride_y                      Stride of the first source image in Y dimension (in bytes)
+ * @param[in]  output_step_y                        output_stride_y * number of elements along Y processed per workitem(in bytes)
+ * @param[in]  output_offset_first_element_in_bytes The offset of the first element in the first source image
+ * @param[in]  threshold_value                      Threshold value.
+ *
+ */
+__kernel void fast_corners(
+    IMAGE_DECLARATION(input),
+    IMAGE_DECLARATION(output),
+    float threshold_value)
+{
+    Image in  = CONVERT_TO_IMAGE_STRUCT(input);
+    Image out = CONVERT_TO_IMAGE_STRUCT(output);
+
+    const unsigned char threshold = (uchar)threshold_value;
+
+    unsigned int bright_mask = 0;
+    unsigned int dark_mask   = 0;
+
+    unsigned char isCorner = 0;
+
+    unsigned char p        = *in.ptr;
+    unsigned char p_bright = add_sat(p, threshold);
+    unsigned char p_dark   = sub_sat(p, threshold);
+
+    LOAD_AND_SET_MASK(in.ptr, 0, input_stride_y, p_dark, p_bright, dark_mask, bright_mask)
+    LOAD_AND_SET_MASK(in.ptr, 4, input_stride_y, p_dark, p_bright, dark_mask, bright_mask)
+    LOAD_AND_SET_MASK(in.ptr, 8, input_stride_y, p_dark, p_bright, dark_mask, bright_mask)
+    LOAD_AND_SET_MASK(in.ptr, 12, input_stride_y, p_dark, p_bright, dark_mask, bright_mask)
+
+    if(((bright_mask | dark_mask) & 0x1111) == 0)
+    {
+        *out.ptr = 0;
+        return;
+    }
+
+    LOAD_AND_SET_MASK(in.ptr, 1, input_stride_y, p_dark, p_bright, dark_mask, bright_mask)
+    LOAD_AND_SET_MASK(in.ptr, 2, input_stride_y, p_dark, p_bright, dark_mask, bright_mask)
+    LOAD_AND_SET_MASK(in.ptr, 3, input_stride_y, p_dark, p_bright, dark_mask, bright_mask)
+    LOAD_AND_SET_MASK(in.ptr, 5, input_stride_y, p_dark, p_bright, dark_mask, bright_mask)
+    LOAD_AND_SET_MASK(in.ptr, 6, input_stride_y, p_dark, p_bright, dark_mask, bright_mask)
+    LOAD_AND_SET_MASK(in.ptr, 7, input_stride_y, p_dark, p_bright, dark_mask, bright_mask)
+    LOAD_AND_SET_MASK(in.ptr, 9, input_stride_y, p_dark, p_bright, dark_mask, bright_mask)
+    LOAD_AND_SET_MASK(in.ptr, 10, input_stride_y, p_dark, p_bright, dark_mask, bright_mask)
+    LOAD_AND_SET_MASK(in.ptr, 11, input_stride_y, p_dark, p_bright, dark_mask, bright_mask)
+    LOAD_AND_SET_MASK(in.ptr, 13, input_stride_y, p_dark, p_bright, dark_mask, bright_mask)
+    LOAD_AND_SET_MASK(in.ptr, 14, input_stride_y, p_dark, p_bright, dark_mask, bright_mask)
+    LOAD_AND_SET_MASK(in.ptr, 15, input_stride_y, p_dark, p_bright, dark_mask, bright_mask)
+
+    bright_mask |= (bright_mask << 16);
+    dark_mask |= (dark_mask << 16);
+
+    CHECK_CORNER(bright_mask, dark_mask, isCorner)
+
+    if(!isCorner)
+    {
+        *out.ptr = 0;
+        return;
+    }
+
+#ifndef USE_MAXSUPPRESSION
+    *out.ptr = 1;
+#else
+
+    *out.ptr = compute_strength(p, in.ptr, input_stride_y, threshold);
+#endif
+}
+
+/** Copy result to Keypoint buffer and count number of corners
+ *
+ * @param[in]  input_ptr                           Pointer to the image with calculated strenghs. Supported data types: U8
+ * @param[in]  input_stride_x                      Stride of the first source image in X dimension (in bytes)
+ * @param[in]  input_step_x                        input_stride_x * number of elements along X processed per workitem(in bytes)
+ * @param[in]  input_stride_y                      Stride of the first source image in Y dimension (in bytes)
+ * @param[in]  input_step_y                        input_stride_y * number of elements along Y processed per workitem(in bytes)
+ * @param[in]  input_offset_first_element_in_bytes The offset of the first element in the first source image
+ * @param[in]  max_num_points                      The maximum number of keypoints the array can hold
+ * @param[out] offset                              The number of skipped pixels in x dimension
+ * @param[out] num_of_points                       Number of points found
+ * @param[out] out                                 The keypoints found
+ *
+ */
+__kernel void copy_to_keypoint(
+    IMAGE_DECLARATION(input),
+    uint     max_num_points,
+    uint     offset,
+    __global uint *num_of_points,
+    __global Keypoint *out)
+{
+#ifndef UPDATE_NUMBER
+    if(*num_of_points >= max_num_points)
+    {
+        return;
+    }
+#endif
+
+    Image in = CONVERT_TO_IMAGE_STRUCT(input);
+
+    uchar value = *in.ptr;
+
+    if(value > 0)
+    {
+        int id = atomic_inc(num_of_points);
+        if(id < max_num_points)
+        {
+            out[id].strength        = value;
+            out[id].x               = get_global_id(0) + offset;
+            out[id].y               = get_global_id(1) + offset;
+            out[id].tracking_status = 1;
+        }
+    }
+}