diff options
Diffstat (limited to 'src/core/CL/cl_kernels/fast_corners.cl')
| -rw-r--r-- | src/core/CL/cl_kernels/fast_corners.cl | 262 |
1 files changed, 0 insertions, 262 deletions
diff --git a/src/core/CL/cl_kernels/fast_corners.cl b/src/core/CL/cl_kernels/fast_corners.cl deleted file mode 100644 index 76b35b931b..0000000000 --- a/src/core/CL/cl_kernels/fast_corners.cl +++ /dev/null @@ -1,262 +0,0 @@ -/* - * 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; - } - } -} |