diff options
| author | Abe Mbise <abe.mbise@arm.com> | 2017-12-19 13:00:58 +0000 |
|---|---|---|
| committer | Anthony Barbier <anthony.barbier@arm.com> | 2018-11-02 16:46:07 +0000 |
| commit | 25a340f84cdecb2c15fde251e59cc1af48648031 (patch) | |
| tree | fc64ba0ed6d814bd91bd88124ba52596dc99d8e1 /src | |
| parent | 40df1e9b770393a77afc2ebb94ff7e8f6f8696ed (diff) | |
| download | ComputeLibrary-25a340f84cdecb2c15fde251e59cc1af48648031.tar.gz | |
COMPMID-578: Implement FAST corners for CL/NEON
Change-Id: Ifa74e2bf05546de9a49aa185e22fba50438d8ad6
Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/113946
Tested-by: Jenkins <bsgcomp@arm.com>
Reviewed-by: Pablo Tello <pablo.tello@arm.com>
Diffstat (limited to 'src')
| -rw-r--r-- | src/core/CL/cl_kernels/fast_corners.cl | 5 | ||||
| -rw-r--r-- | src/core/CL/kernels/CLFastCornersKernel.cpp | 4 | ||||
| -rw-r--r-- | src/core/NEON/kernels/NEFillArrayKernel.cpp | 5 | ||||
| -rw-r--r-- | src/runtime/CL/functions/CLFastCorners.cpp | 8 |
4 files changed, 14 insertions, 8 deletions
diff --git a/src/core/CL/cl_kernels/fast_corners.cl b/src/core/CL/cl_kernels/fast_corners.cl index 3e1929c637..76b35b931b 100644 --- a/src/core/CL/cl_kernels/fast_corners.cl +++ b/src/core/CL/cl_kernels/fast_corners.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016, 2017 ARM Limited. + * Copyright (c) 2016-2018 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -254,6 +254,9 @@ __kernel void copy_to_keypoint( 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; } } } diff --git a/src/core/CL/kernels/CLFastCornersKernel.cpp b/src/core/CL/kernels/CLFastCornersKernel.cpp index 1d4d776730..616e41b5fc 100644 --- a/src/core/CL/kernels/CLFastCornersKernel.cpp +++ b/src/core/CL/kernels/CLFastCornersKernel.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016, 2017 ARM Limited. + * Copyright (c) 2016-2018 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -138,7 +138,7 @@ void CLCopyToArrayKernel::configure(const ICLImage *input, bool update_number, I // Set static kernel arguments unsigned int idx = num_arguments_per_2D_tensor(); // Skip the input and output parameters - _kernel.setArg<unsigned int>(idx++, corners->max_num_values()); + _kernel.setArg<unsigned int>(idx++, _corners->max_num_values()); _kernel.setArg<cl_uint>(idx++, offset); _kernel.setArg(idx++, *_num_buffer); _kernel.setArg(idx++, _corners->cl_buffer()); diff --git a/src/core/NEON/kernels/NEFillArrayKernel.cpp b/src/core/NEON/kernels/NEFillArrayKernel.cpp index 5a2e1a0aa4..c6c6c452b7 100644 --- a/src/core/NEON/kernels/NEFillArrayKernel.cpp +++ b/src/core/NEON/kernels/NEFillArrayKernel.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016, 2017 ARM Limited. + * Copyright (c) 2016-2018 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -81,6 +81,9 @@ void NEFillArrayKernel::run(const Window &window, const ThreadInfo &info) p.y = id.y(); p.strength = value; p.tracking_status = 1; + p.scale = 0.f; + p.orientation = 0.f; + p.error = 0.f; if(!_output->push_back(p)) { diff --git a/src/runtime/CL/functions/CLFastCorners.cpp b/src/runtime/CL/functions/CLFastCorners.cpp index 7a0dd09fbe..d6cda91cea 100644 --- a/src/runtime/CL/functions/CLFastCorners.cpp +++ b/src/runtime/CL/functions/CLFastCorners.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016, 2017 ARM Limited. + * Copyright (c) 2016-2018 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -52,7 +52,7 @@ CLFastCorners::CLFastCorners(std::shared_ptr<IMemoryManager> memory_manager) { } -void CLFastCorners::configure(const ICLImage *input, float threshold, bool nonmax_suppression, CLKeyPointArray *const corners, +void CLFastCorners::configure(const ICLImage *input, float threshold, bool nonmax_suppression, ICLKeyPointArray *corners, unsigned int *num_corners, BorderMode border_mode, uint8_t constant_border_value) { ARM_COMPUTE_ERROR_ON_TENSOR_NOT_2D(input); @@ -76,7 +76,7 @@ void CLFastCorners::configure(const ICLImage *input, float threshold, bool nonma if(!_non_max) { - _copy_array_kernel.configure(&_output, update_number, corners, &_num_buffer); + _copy_array_kernel.configure(&_output, update_number, _corners, &_num_buffer); } else { @@ -84,7 +84,7 @@ void CLFastCorners::configure(const ICLImage *input, float threshold, bool nonma _memory_group.manage(&_suppr); _suppr_func.configure(&_output, &_suppr, border_mode); - _copy_array_kernel.configure(&_suppr, update_number, corners, &_num_buffer); + _copy_array_kernel.configure(&_suppr, update_number, _corners, &_num_buffer); _suppr.allocator()->allocate(); } |