src/core/CL/cl_kernels/fast_corners.cl

/*
 * Copyright (c) 2016-2018 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;
    }

#ifdef USE_MAXSUPPRESSION
    *out.ptr = compute_strength(p, in.ptr, input_stride_y, threshold);
#else  /* USE_MAXSUPPRESSION */
    *out.ptr = 1;
#endif /* USE_MAXSUPPRESSION */
}

/** 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 /* UPDATE_NUMBER */

    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;
            out[id].scale           = 0.f;
            out[id].orientation     = 0.f;
            out[id].error           = 0.f;
        }
    }
}