fast_corners.cl

Go to the documentation of this file.

/*
 * 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
};

#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;                                   \
    }

#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;
}

__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 */
}

__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;
        }
    }
}