Remove (NE/CL)YoloLayer support

YOLO layer is too specialized and specific to a single model type. Can
be decomposed using split, activation and concatenate layers

Partially Resolves: COMPMID-3996

Signed-off-by: Georgios Pinitas <georgios.pinitas@arm.com>
Change-Id: I3cde88f8d4cc7d8c70ce1bb3b32b00f8d09bdca2
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/4678
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Michele Di Giorgio <michele.digiorgio@arm.com>

6 files changed, 0 insertions, 603 deletions

diff --git a/src/core/CL/CLKernels.h b/src/core/CL/CLKernels.h
index a9654ecc1f..42fe79e91d 100644
--- a/src/core/CL/CLKernels.h
+++ b/src/core/CL/CLKernels.h
@@ -154,7 +154,6 @@
 #include "src/core/CL/kernels/CLWinogradFilterTransformKernel.h"
 #include "src/core/CL/kernels/CLWinogradInputTransformKernel.h"
 #include "src/core/CL/kernels/CLWinogradOutputTransformKernel.h"
-#include "src/core/CL/kernels/CLYOLOLayerKernel.h"
 #include "src/core/CL/kernels/ICLDepthwiseConvolutionLayer3x3Kernel.h"
 
 #endif /* ARM_COMPUTE_CLKERNELS_H */
diff --git a/src/core/CL/kernels/CLYOLOLayerKernel.cpp b/src/core/CL/kernels/CLYOLOLayerKernel.cpp
deleted file mode 100644
index e12d1e7a65..0000000000
--- a/src/core/CL/kernels/CLYOLOLayerKernel.cpp
+++ /dev/null
@@ -1,187 +0,0 @@
-/*
- * Copyright (c) 2018-2020 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 "src/core/CL/kernels/CLYOLOLayerKernel.h"
-
-#include "arm_compute/core/CL/CLHelpers.h"
-#include "arm_compute/core/CL/CLHelpers.h"
-#include "arm_compute/core/CL/CLKernelLibrary.h"
-#include "arm_compute/core/CL/ICLTensor.h"
-#include "arm_compute/core/Helpers.h"
-#include "arm_compute/core/IAccessWindow.h"
-#include "arm_compute/core/TensorInfo.h"
-#include "arm_compute/core/Types.h"
-#include "arm_compute/core/Utils.h"
-#include "arm_compute/core/Window.h"
-#include "src/core/CL/CLValidate.h"
-#include "src/core/helpers/AutoConfiguration.h"
-#include "src/core/helpers/WindowHelpers.h"
-
-#include "support/StringSupport.h"
-
-namespace arm_compute
-{
-namespace
-{
-Status validate_arguments(const ITensorInfo *input, const ITensorInfo *output, const ActivationLayerInfo &act_info, int32_t num_classes)
-{
-    ARM_COMPUTE_RETURN_ERROR_ON_F16_UNSUPPORTED(input);
-    ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::F16, DataType::F32);
-    ARM_COMPUTE_RETURN_ERROR_ON(input->data_layout() == DataLayout::UNKNOWN);
-    ARM_COMPUTE_RETURN_ERROR_ON(act_info.activation() != ActivationLayerInfo::ActivationFunction::LOGISTIC);
-
-    const unsigned int channel_idx = get_data_layout_dimension_index(input->data_layout(), DataLayoutDimension::CHANNEL);
-    ARM_COMPUTE_RETURN_ERROR_ON(num_classes <= 0);
-    ARM_COMPUTE_RETURN_ERROR_ON((input->dimension(channel_idx) % (num_classes + 5)) != 0);
-
-    // Checks performed when output is configured
-    if((output != nullptr) && (output->total_size() != 0))
-    {
-        ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_SHAPES(input, output);
-        ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_TYPES(input, output);
-    }
-
-    return Status{};
-}
-
-std::pair<Status, Window> validate_and_configure_window(ITensorInfo *input, ITensorInfo *output)
-{
-    if(output != nullptr)
-    {
-        ARM_COMPUTE_ERROR_ON_NULLPTR(input, output);
-
-        // Output auto inizialitation if not yet initialized
-        auto_init_if_empty(*output, *input);
-    }
-
-    const bool         is_nchw                           = input->data_layout() == DataLayout::NCHW;
-    const unsigned int num_elems_processed_per_iteration = is_nchw ? 16 / input->element_size() : 1;
-
-    Window win            = calculate_max_window(*input, Steps(num_elems_processed_per_iteration));
-    bool   window_changed = false;
-
-    if(output != nullptr)
-    {
-        AccessWindowHorizontal input_access(input, 0, num_elems_processed_per_iteration);
-        AccessWindowHorizontal output_access(output, 0, num_elems_processed_per_iteration);
-        window_changed = update_window_and_padding(win, input_access, output_access);
-        output_access.set_valid_region(win, input->valid_region());
-    }
-    else
-    {
-        window_changed = update_window_and_padding(win, AccessWindowHorizontal(input, 0, num_elems_processed_per_iteration));
-    }
-
-    Status err = (window_changed) ? ARM_COMPUTE_CREATE_ERROR(ErrorCode::RUNTIME_ERROR, "Insufficient Padding!") : Status{};
-    return std::make_pair(err, win);
-}
-} // namespace
-
-CLYOLOLayerKernel::CLYOLOLayerKernel()
-    : _input(nullptr), _output(nullptr), _run_in_place(false)
-{
-}
-
-void CLYOLOLayerKernel::configure(ICLTensor *input, ICLTensor *output, const ActivationLayerInfo &act_info, int32_t num_classes)
-{
-    configure(CLKernelLibrary::get().get_compile_context(), input, output, act_info, num_classes);
-}
-
-void CLYOLOLayerKernel::configure(const CLCompileContext &compile_context, ICLTensor *input, ICLTensor *output, const ActivationLayerInfo &act_info, int32_t num_classes)
-{
-    ARM_COMPUTE_ERROR_ON_NULLPTR(input);
-
-    _run_in_place = (output == nullptr) || (output == input);
-
-    ARM_COMPUTE_ERROR_THROW_ON(validate_arguments(input->info(), (output != nullptr) ? output->info() : nullptr, act_info, num_classes));
-
-    const bool         is_nchw                           = input->info()->data_layout() == DataLayout::NCHW;
-    const unsigned int num_elems_processed_per_iteration = is_nchw ? 16 / input->info()->element_size() : 1;
-    const DataType     dt                                = input->info()->data_type();
-    float              a_const                           = act_info.a();
-    float              b_const                           = act_info.b();
-
-    // Set build options
-    CLBuildOptions build_opts;
-    build_opts.add_option("-DACTIVATION_TYPE=" + lower_string(string_from_activation_func(act_info.activation())));
-    build_opts.add_option("-DDATA_TYPE=" + get_cl_type_from_data_type(dt));
-    build_opts.add_option("-DVEC_SIZE=" + support::cpp11::to_string(num_elems_processed_per_iteration));
-    build_opts.add_option("-DA_VAL=" + float_to_string_with_full_precision(a_const));
-    build_opts.add_option("-DB_VAL=" + float_to_string_with_full_precision(b_const));
-    build_opts.add_option("-DNUM_CLASSES=" + support::cpp11::to_string(num_classes));
-    build_opts.add_option_if(_run_in_place, "-DIN_PLACE");
-
-    // Create kernel
-    std::string kernel_name = std::string("yolo_layer_") + lower_string(string_from_data_layout(input->info()->data_layout()));
-    _kernel                 = create_kernel(compile_context, kernel_name, build_opts.options());
-
-    // Make sure _kernel is initialized before calling the parent's configure
-    _input  = input;
-    _output = output;
-
-    // Configure kernel window
-    auto win_config = validate_and_configure_window(input->info(), (_run_in_place) ? nullptr : output->info());
-    ARM_COMPUTE_ERROR_THROW_ON(win_config.first);
-    ICLKernel::configure_internal(win_config.second);
-
-    // Set config_id for enabling LWS tuning
-    _config_id = "yolo_layer_";
-    _config_id += lower_string(string_from_data_type(dt));
-    _config_id += "_";
-    _config_id += support::cpp11::to_string(input->info()->dimension(0));
-    _config_id += "_";
-    _config_id += support::cpp11::to_string(input->info()->dimension(1));
-    _config_id += "_";
-    _config_id += lower_string(string_from_data_layout(input->info()->data_layout()));
-}
-
-Status CLYOLOLayerKernel::validate(const ITensorInfo *input, const ITensorInfo *output, const ActivationLayerInfo &act_info, int32_t num_classes)
-{
-    const bool run_in_place = (output == nullptr) || (output == input);
-    ARM_COMPUTE_RETURN_ON_ERROR(validate_arguments(input, output, act_info, num_classes));
-    ARM_COMPUTE_RETURN_ON_ERROR(validate_and_configure_window(input->clone().get(), (run_in_place) ? nullptr : output->clone().get()).first);
-
-    return Status{};
-}
-
-void CLYOLOLayerKernel::run(const Window &window, cl::CommandQueue &queue)
-{
-    ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this);
-    ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW(ICLKernel::window(), window);
-
-    Window collapsed = window.collapse_if_possible(ICLKernel::window(), Window::DimZ);
-    Window slice     = collapsed.first_slice_window_3D();
-
-    do
-    {
-        unsigned int idx = 0;
-        add_3D_tensor_argument(idx, _input, slice);
-        if(!_run_in_place)
-        {
-            add_3D_tensor_argument(idx, _output, slice);
-        }
-        enqueue(queue, *this, slice, lws_hint());
-    }
-    while(collapsed.slide_window_slice_3D(slice));
-}
-} // namespace arm_compute
diff --git a/src/core/CL/kernels/CLYOLOLayerKernel.h b/src/core/CL/kernels/CLYOLOLayerKernel.h
deleted file mode 100644
index 5b1d56e9e5..0000000000
--- a/src/core/CL/kernels/CLYOLOLayerKernel.h
+++ /dev/null
@@ -1,98 +0,0 @@
-/*
- * Copyright (c) 2018-2020 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.
- */
-#ifndef ARM_COMPUTE_CLYOLOLAYERKERNEL_H
-#define ARM_COMPUTE_CLYOLOLAYERKERNEL_H
-
-#include "src/core/CL/ICLKernel.h"
-
-namespace arm_compute
-{
-class ICLTensor;
-
-/** Interface for the YOLO layer kernel that performs partial activation.
- *  For each box, activate only:
- *    - x and y position (channel 0 and 1 of each box)
- *    - objectiveness    (channel 4 of each box)
- *    - classes          (channel 5 to (classes - 5) of each box)
- */
-class CLYOLOLayerKernel : public ICLKernel
-{
-public:
-    /** Default constructor */
-    CLYOLOLayerKernel();
-    /** Prevent instances of this class from being copied (As this class contains pointers) */
-    CLYOLOLayerKernel(const CLYOLOLayerKernel &) = delete;
-    /** Prevent instances of this class from being copied (As this class contains pointers) */
-    CLYOLOLayerKernel &operator=(const CLYOLOLayerKernel &) = delete;
-    /** Allow instances of this class to be moved */
-    CLYOLOLayerKernel(CLYOLOLayerKernel &&) = default;
-    /** Allow instances of this class to be moved */
-    CLYOLOLayerKernel &operator=(CLYOLOLayerKernel &&) = default;
-    /** Default destructor */
-    ~CLYOLOLayerKernel() = default;
-    /** Set the input and output tensor.
-     *
-     * @note If the output tensor is a nullptr, the activation function will be performed in-place
-     *
-     * @param[in, out] input       Source tensor. In case of @p output tensor = nullptr, this tensor will store the result
-     *                             of the activation function. Data types supported: F16/F32.
-     * @param[out]     output      Destination tensor. Data type supported: same as @p input
-     * @param[in]      act_info    Activation layer information.
-     * @param[in]      num_classes Number of classes to activate (must be submultiple of @p input channels)
-     */
-    void configure(ICLTensor *input, ICLTensor *output, const ActivationLayerInfo &act_info, int32_t num_classes);
-    /** Set the input and output tensor.
-     *
-     * @note If the output tensor is a nullptr, the activation function will be performed in-place
-     *
-     * @param[in]      compile_context The compile context to be used.
-     * @param[in, out] input           Source tensor. In case of @p output tensor = nullptr, this tensor will store the result
-     *                                 of the activation function. Data types supported: F16/F32.
-     * @param[out]     output          Destination tensor. Data type supported: same as @p input
-     * @param[in]      act_info        Activation layer information.
-     * @param[in]      num_classes     Number of classes to activate (must be submultiple of @p input channels)
-     */
-    void configure(const CLCompileContext &compile_context, ICLTensor *input, ICLTensor *output, const ActivationLayerInfo &act_info, int32_t num_classes);
-    /** Static function to check if given info will lead to a valid configuration of @ref CLYOLOLayerKernel
-     *
-     * @param[in] input       Source tensor info. In case of @p output tensor info = nullptr, this tensor will store the result
-     *                        of the activation function. Data types supported: F16/F32.
-     * @param[in] output      Destination tensor info. Data type supported: same as @p input
-     * @param[in] act_info    Activation layer information.
-     * @param[in] num_classes Number of classes to activate (must be submultiple of @p input channels)
-     *
-     * @return a status
-     */
-    static Status validate(const ITensorInfo *input, const ITensorInfo *output, const ActivationLayerInfo &act_info, int32_t num_classes);
-
-    // Inherited methods overridden:
-    void run(const Window &window, cl::CommandQueue &queue) override;
-
-private:
-    ICLTensor *_input;
-    ICLTensor *_output;
-    bool       _run_in_place;
-};
-} // namespace arm_compute
-#endif /*ARM_COMPUTE_CLYOLOLAYERKERNEL_H */
diff --git a/src/core/NEON/NEKernels.h b/src/core/NEON/NEKernels.h
index 55aa514f36..90c7df3c1c 100644
--- a/src/core/NEON/NEKernels.h
+++ b/src/core/NEON/NEKernels.h
@@ -145,6 +145,5 @@
 #include "src/core/NEON/kernels/NEWeightsReshapeKernel.h"
 #include "src/core/NEON/kernels/NEWidthConcatenateLayerKernel.h"
 #include "src/core/NEON/kernels/NEWinogradConvolutionLayerKernel.h"
-#include "src/core/NEON/kernels/NEYOLOLayerKernel.h"
 
 #endif /* ARM_COMPUTE_NEKERNELS_H */
diff --git a/src/core/NEON/kernels/NEYOLOLayerKernel.cpp b/src/core/NEON/kernels/NEYOLOLayerKernel.cpp
deleted file mode 100644
index 33bcc20d39..0000000000
--- a/src/core/NEON/kernels/NEYOLOLayerKernel.cpp
+++ /dev/null
@@ -1,210 +0,0 @@
-/*
- * Copyright (c) 2018-2020 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 "src/core/NEON/kernels/NEYOLOLayerKernel.h"
-
-#include "arm_compute/core/Helpers.h"
-#include "arm_compute/core/ITensor.h"
-#include "arm_compute/core/TensorInfo.h"
-#include "arm_compute/core/Utils.h"
-#include "arm_compute/core/Validate.h"
-#include "arm_compute/core/Window.h"
-#include "src/core/CPP/Validate.h"
-#include "src/core/NEON/NEAsymm.h"
-#include "src/core/NEON/NEFixedPoint.h"
-#include "src/core/NEON/NEMath.h"
-#include "src/core/helpers/AutoConfiguration.h"
-#include "src/core/helpers/WindowHelpers.h"
-
-#include "src/core/NEON/kernels/detail/NEActivationFunctionDetail.h"
-
-#include <arm_neon.h>
-
-namespace arm_compute
-{
-namespace
-{
-Status validate_arguments(const ITensorInfo *input, const ITensorInfo *output, const ActivationLayerInfo &act_info, int32_t num_classes)
-{
-    ARM_COMPUTE_UNUSED(act_info);
-    ARM_COMPUTE_RETURN_ERROR_ON_CPU_F16_UNSUPPORTED(input);
-    ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::F16, DataType::F32);
-    ARM_COMPUTE_RETURN_ERROR_ON(input->data_layout() == DataLayout::UNKNOWN);
-    ARM_COMPUTE_RETURN_ERROR_ON(act_info.activation() != ActivationLayerInfo::ActivationFunction::LOGISTIC);
-
-    const unsigned int channel_idx = get_data_layout_dimension_index(input->data_layout(), DataLayoutDimension::CHANNEL);
-    ARM_COMPUTE_RETURN_ERROR_ON(num_classes <= 0);
-    ARM_COMPUTE_RETURN_ERROR_ON((input->dimension(channel_idx) % (num_classes + 5)) != 0);
-
-    // Checks performed when output is configured
-    if((output != nullptr) && (output->total_size() != 0))
-    {
-        ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_SHAPES(input, output);
-        ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_TYPES(input, output);
-    }
-
-    return Status{};
-}
-} // namespace
-
-NEYOLOLayerKernel::NEYOLOLayerKernel()
-    : _func(nullptr), _input(nullptr), _output(nullptr), _act_info(), _num_classes()
-{
-}
-
-template <typename T, int S>
-void NEYOLOLayerKernel::yolo_layer_nchw(const Window &window)
-{
-    const auto window_start_x = static_cast<int>(window.x().start());
-    const auto window_end_x   = static_cast<int>(window.x().end());
-    const int  window_step_x  = S;
-
-    Window win{ window };
-    win.set(Window::DimX, Window::Dimension(0, 1, 1));
-    Iterator input(_input, win);
-    Iterator output(_output, win);
-
-    execute_window_loop(win, [&](const Coordinates & id)
-    {
-        const auto input_ptr  = reinterpret_cast<const T *>(input.ptr());
-        const auto output_ptr = reinterpret_cast<T *>(output.ptr());
-        int        x          = window_start_x;
-        const int  box_ch_id  = id.z() % (_num_classes + 5);
-        const bool activate   = box_ch_id != 2 && box_ch_id != 3;
-
-        for(; x <= (window_end_x - window_step_x); x += window_step_x)
-        {
-            auto res = wrapper::vloadq(input_ptr + x);
-
-            // Perform activation
-            if(activate)
-            {
-                auto activation = detail::logistic<T, S>(_act_info);
-                activation(res);
-            }
-
-            // Store results
-            wrapper::vstore(output_ptr + x, res);
-        }
-
-        // Compute left-over elements
-        for(; x < window_end_x; ++x)
-        {
-            auto res = *(input_ptr + x);
-
-            // Perform activation
-            if(activate)
-            {
-                res = 1.f / (1.f + std::exp(-res));
-            }
-
-            *(output_ptr + x) = res;
-        }
-    },
-    input, output);
-}
-
-template <typename T>
-void NEYOLOLayerKernel::yolo_layer_nhwc(const Window &window)
-{
-    Iterator input(_input, window);
-    Iterator output(_output, window);
-
-    execute_window_loop(window, [&](const Coordinates & id)
-    {
-        auto res = *(reinterpret_cast<T *>(input.ptr()));
-
-        const int  box_ch_id = id.x() % (_num_classes + 5);
-        const bool activate  = box_ch_id != 2 && box_ch_id != 3;
-
-        // Perform activation
-        if(activate)
-        {
-            res = 1.f / (1.f + std::exp(-res));
-        }
-
-        // Store result
-        *(reinterpret_cast<T *>(output.ptr())) = res;
-    },
-    input, output);
-}
-
-void NEYOLOLayerKernel::configure(ITensor *input, ITensor *output, const ActivationLayerInfo &act_info, int32_t num_classes)
-{
-    ARM_COMPUTE_ERROR_ON_NULLPTR(input);
-    ARM_COMPUTE_ERROR_THROW_ON(validate_arguments(input->info(), (output != nullptr) ? output->info() : nullptr, act_info, num_classes));
-
-    _input       = input;
-    _output      = output;
-    _act_info    = act_info;
-    _num_classes = num_classes;
-
-    switch(_input->info()->data_type())
-    {
-#ifdef __ARM_FEATURE_FP16_VECTOR_ARITHMETIC
-        case DataType::F16:
-            _func = (_input->info()->data_layout() == DataLayout::NHWC) ? &NEYOLOLayerKernel::yolo_layer_nhwc<float16_t> : &NEYOLOLayerKernel::yolo_layer_nchw<float16_t, 8>;
-            break;
-#endif // __ARM_FEATURE_FP16_VECTOR_ARITHMETIC
-        case DataType::F32:
-            _func = (_input->info()->data_layout() == DataLayout::NHWC) ? &NEYOLOLayerKernel::yolo_layer_nhwc<float> : &NEYOLOLayerKernel::yolo_layer_nchw<float, 4>;
-            break;
-        default:
-            ARM_COMPUTE_ERROR("Element size not supported");
-            break;
-    }
-
-    Window win = calculate_max_window(*input->info(), Steps());
-
-    // Configure kernel window
-    if(output != nullptr)
-    {
-        // Output auto inizialitation if not yet initialized
-        auto_init_if_empty(*output->info(), *input->info());
-
-        Coordinates coord;
-        coord.set_num_dimensions(output->info()->num_dimensions());
-
-        output->info()->set_valid_region(ValidRegion(coord, output->info()->tensor_shape()));
-    }
-
-    ICPPKernel::configure(win);
-}
-
-Status NEYOLOLayerKernel::validate(const ITensorInfo *input, const ITensorInfo *output, const ActivationLayerInfo &act_info, int32_t num_classes)
-{
-    ARM_COMPUTE_RETURN_ON_ERROR(validate_arguments(input, output, act_info, num_classes));
-
-    return Status{};
-}
-
-void NEYOLOLayerKernel::run(const Window &window, const ThreadInfo &info)
-{
-    ARM_COMPUTE_UNUSED(info);
-    ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this);
-    ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW(IKernel::window(), window);
-    ARM_COMPUTE_ERROR_ON(_func == nullptr);
-
-    (this->*_func)(window);
-}
-} // namespace arm_compute
diff --git a/src/core/NEON/kernels/NEYOLOLayerKernel.h b/src/core/NEON/kernels/NEYOLOLayerKernel.h
deleted file mode 100644
index 806cf9cc09..0000000000
--- a/src/core/NEON/kernels/NEYOLOLayerKernel.h
+++ /dev/null
@@ -1,106 +0,0 @@
-/*
- * Copyright (c) 2018-2020 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.
- */
-#ifndef ARM_COMPUTE_NEYOLOLAYERKERNEL_H
-#define ARM_COMPUTE_NEYOLOLAYERKERNEL_H
-
-#include "src/core/NEON/INEKernel.h"
-
-namespace arm_compute
-{
-class ITensor;
-
-/** Interface for the YOLO layer kernel. */
-class NEYOLOLayerKernel : public INEKernel
-{
-public:
-    const char *name() const override
-    {
-        return "NEYOLOLayerKernel";
-    }
-    /** Constructor */
-    NEYOLOLayerKernel();
-    /** Prevent instances of this class from being copied (As this class contains pointers) */
-    NEYOLOLayerKernel(const NEYOLOLayerKernel &) = delete;
-    /** Default move constructor */
-    NEYOLOLayerKernel(NEYOLOLayerKernel &&) = default;
-    /** Prevent instances of this class from being copied (As this class contains pointers) */
-    NEYOLOLayerKernel &operator=(const NEYOLOLayerKernel &) = delete;
-    /** Default move assignment operator */
-    NEYOLOLayerKernel &operator=(NEYOLOLayerKernel &&) = default;
-    /** Default destructor */
-    ~NEYOLOLayerKernel() = default;
-    /** Set the input and output tensor.
-     *
-     * @note If the output tensor is a nullptr or is equal to the input, the activation function will be performed in-place
-     *
-     * @param[in, out] input       Source tensor. In case of @p output tensor = nullptr, this tensor will store the result
-     *                             of the activation function. Data types supported: F16/F32.
-     * @param[out]     output      Destination tensor. Data type supported: same as @p input
-     * @param[in]      act_info    Activation layer parameters.
-     * @param[in]      num_classes Number of classes to activate (must be submultiple of @p input channels)
-     */
-    void configure(ITensor *input, ITensor *output, const ActivationLayerInfo &act_info, int32_t num_classes);
-    /** Static function to check if given info will lead to a valid configuration of @ref NEYOLOLayerKernel
-     *
-     * @param[in] input       Source tensor info. In case of @p output tensor info = nullptr, this tensor will store the result
-     *                        of the activation function. Data types supported: F16/F32.
-     * @param[in] output      Destination tensor info. Data type supported: same as @p input
-     * @param[in] act_info    Activation layer information.
-     * @param[in] num_classes Number of classes to activate (must be submultiple of @p input channels)
-     *
-     * @return a status
-     */
-    static Status validate(const ITensorInfo *input, const ITensorInfo *output, const ActivationLayerInfo &act_info, int32_t num_classes);
-
-    // Inherited methods overridden:
-    void run(const Window &window, const ThreadInfo &info) override;
-
-private:
-    /** Function to run YOLO layer
-     *
-     * @param[in] window Region on which to execute the kernel.
-     */
-    template <typename T, int S>
-    void yolo_layer_nchw(const Window &window);
-    /** Function to run YOLO layer on tensors with NHWC format
-     *
-     * @param[in] window Region on which to execute the kernel.
-     */
-    template <typename T>
-    void yolo_layer_nhwc(const Window &window);
-    /** Common signature for all the yolo layer functions
-     *
-     * @param[in] window Region on which to execute the kernel.
-     */
-    using YOLOFunctionPtr = void (NEYOLOLayerKernel::*)(const Window &window);
-
-private:
-    YOLOFunctionPtr     _func;
-    ITensor            *_input;
-    ITensor            *_output;
-    ActivationLayerInfo _act_info;
-    int32_t             _num_classes;
-};
-} // namespace arm_compute
-#endif /*ARM_COMPUTE_NEYOLOLAYERKERNEL_H */