Remove support for (NE/CL)LocallyConnectedLayer

Remove out-of-date and unmaintained LocallyConnectedLayer for both NEON
and OpenCL.

Resolves: COMPMID-3924

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

15 files changed, 8 insertions, 1446 deletions

diff --git a/Android.bp b/Android.bp
index e9b3588a83..bf68dd4535 100644
--- a/Android.bp
+++ b/Android.bp
@@ -156,7 +156,6 @@ cc_library_static {
         "src/core/CL/kernels/CLIntegralImageKernel.cpp",
         "src/core/CL/kernels/CLL2NormalizeLayerKernel.cpp",
         "src/core/CL/kernels/CLLKTrackerKernel.cpp",
-        "src/core/CL/kernels/CLLocallyConnectedMatrixMultiplyKernel.cpp",
         "src/core/CL/kernels/CLMagnitudePhaseKernel.cpp",
         "src/core/CL/kernels/CLMaxUnpoolingLayerKernel.cpp",
         "src/core/CL/kernels/CLMeanStdDevKernel.cpp",
@@ -305,7 +304,6 @@ cc_library_static {
         "src/core/NEON/kernels/NEIntegralImageKernel.cpp",
         "src/core/NEON/kernels/NEL2NormalizeLayerKernel.cpp",
         "src/core/NEON/kernels/NELKTrackerKernel.cpp",
-        "src/core/NEON/kernels/NELocallyConnectedMatrixMultiplyKernel.cpp",
         "src/core/NEON/kernels/NELogicalKernel.cpp",
         "src/core/NEON/kernels/NEMagnitudePhaseKernel.cpp",
         "src/core/NEON/kernels/NEMaxUnpoolingLayerKernel.cpp",
@@ -529,7 +527,6 @@ cc_library_static {
         "src/runtime/CL/functions/CLLSTMLayerQuantized.cpp",
         "src/runtime/CL/functions/CLLaplacianPyramid.cpp",
         "src/runtime/CL/functions/CLLaplacianReconstruct.cpp",
-        "src/runtime/CL/functions/CLLocallyConnectedLayer.cpp",
         "src/runtime/CL/functions/CLLogicalAnd.cpp",
         "src/runtime/CL/functions/CLLogicalNot.cpp",
         "src/runtime/CL/functions/CLLogicalOr.cpp",
@@ -700,7 +697,6 @@ cc_library_static {
         "src/runtime/NEON/functions/NELSTMLayerQuantized.cpp",
         "src/runtime/NEON/functions/NELaplacianPyramid.cpp",
         "src/runtime/NEON/functions/NELaplacianReconstruct.cpp",
-        "src/runtime/NEON/functions/NELocallyConnectedLayer.cpp",
         "src/runtime/NEON/functions/NELogical.cpp",
         "src/runtime/NEON/functions/NEMagnitude.cpp",
         "src/runtime/NEON/functions/NEMaxUnpoolingLayer.cpp",
diff --git a/arm_compute/runtime/CL/CLFunctions.h b/arm_compute/runtime/CL/CLFunctions.h
index 4e32831bc6..ff686c59d5 100644
--- a/arm_compute/runtime/CL/CLFunctions.h
+++ b/arm_compute/runtime/CL/CLFunctions.h
@@ -98,7 +98,6 @@
 #include "arm_compute/runtime/CL/functions/CLLSTMLayerQuantized.h"
 #include "arm_compute/runtime/CL/functions/CLLaplacianPyramid.h"
 #include "arm_compute/runtime/CL/functions/CLLaplacianReconstruct.h"
-#include "arm_compute/runtime/CL/functions/CLLocallyConnectedLayer.h"
 #include "arm_compute/runtime/CL/functions/CLLogicalAnd.h"
 #include "arm_compute/runtime/CL/functions/CLLogicalNot.h"
 #include "arm_compute/runtime/CL/functions/CLLogicalOr.h"
diff --git a/arm_compute/runtime/CL/functions/CLLocallyConnectedLayer.h b/arm_compute/runtime/CL/functions/CLLocallyConnectedLayer.h
deleted file mode 100644
index 3bbf9f2c30..0000000000
--- a/arm_compute/runtime/CL/functions/CLLocallyConnectedLayer.h
+++ /dev/null
@@ -1,125 +0,0 @@
-/*
- * Copyright (c) 2017-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_CLLOCALLYCONNECTEDLAYER_H
-#define ARM_COMPUTE_CLLOCALLYCONNECTEDLAYER_H
-
-#include "arm_compute/runtime/IFunction.h"
-
-#include "arm_compute/core/Types.h"
-#include "arm_compute/runtime/CL/CLTensor.h"
-#include "arm_compute/runtime/IMemoryManager.h"
-#include "arm_compute/runtime/MemoryGroup.h"
-
-#include <memory>
-
-namespace arm_compute
-{
-class CLCompileContext;
-class CLCol2ImKernel;
-class CLIm2ColKernel;
-class CLWeightsReshapeKernel;
-class CLLocallyConnectedMatrixMultiplyKernel;
-class ICLTensor;
-class ITensorInfo;
-
-/** Basic function to compute the locally connected layer. This function calls the following OpenCL kernels:
- *
- * -# @ref CLWeightsReshapeKernel (executed only once for each configuration)
- * -# @ref CLIm2ColKernel
- * -# @ref CLLocallyConnectedMatrixMultiplyKernel
- * -# @ref CLCol2ImKernel
- */
-class CLLocallyConnectedLayer : public IFunction
-{
-public:
-    /** Default constructor */
-    CLLocallyConnectedLayer(std::shared_ptr<IMemoryManager> memory_manager = nullptr);
-    /** Prevent instances of this class from being copied (As this class contains pointers) */
-    CLLocallyConnectedLayer(const CLLocallyConnectedLayer &) = delete;
-    /** Default move constructor */
-    CLLocallyConnectedLayer(CLLocallyConnectedLayer &&) = default;
-    /** Prevent instances of this class from being copied (As this class contains pointers) */
-    CLLocallyConnectedLayer &operator=(const CLLocallyConnectedLayer &) = delete;
-    /** Default move assignment operator */
-    CLLocallyConnectedLayer &operator=(CLLocallyConnectedLayer &&) = default;
-    /** Set the input and output tensors.
-     *
-     * @param[in]  input     Source tensor. 3 lower dimensions represent a single input [width, height, IFM],
-     *                       while every optional dimension from 4 and above represent a batch of inputs.
-     *                       Data types supported: F32.
-     * @param[in]  weights   Weights tensor. Weights are 5D tensor with dimensions [kernel_x, kernel_y, IFM, OFM, num_patches]. Data type supported:Same as @p input.
-     * @param[in]  biases    Biases tensor. Shared biases supported. Biases are 2D tensor with dimensions [OFM, num_patches]. Data type supported:Same as @p input.
-     * @param[out] output    Destination tensor. 3 lower dimensions represent a single output [width, height, OFM], while the rest represent batch of outputs.
-     *                       Data types supported: Same as @p input.
-     * @param[in]  conv_info Contains padding and stride information described in @ref PadStrideInfo.
-     */
-    ARM_COMPUTE_DEPRECATED_REL(20.11)
-    void configure(const ICLTensor *input, const ICLTensor *weights, const ICLTensor *biases, ICLTensor *output, const PadStrideInfo &conv_info);
-    /** Set the input and output tensors.
-     *
-     * @param[in]  compile_context The compile context to be used.
-     * @param[in]  input           Source tensor. 3 lower dimensions represent a single input [width, height, IFM],
-     *                             while every optional dimension from 4 and above represent a batch of inputs.
-     *                             Data types supported: F32.
-     * @param[in]  weights         Weights tensor. Weights are 5D tensor with dimensions [kernel_x, kernel_y, IFM, OFM, num_patches]. Data type supported:Same as @p input.
-     * @param[in]  biases          Biases tensor. Shared biases supported. Biases are 2D tensor with dimensions [OFM, num_patches]. Data type supported:Same as @p input.
-     * @param[out] output          Destination tensor. 3 lower dimensions represent a single output [width, height, OFM], while the rest represent batch of outputs.
-     *                             Data types supported: Same as @p input.
-     * @param[in]  conv_info       Contains padding and stride information described in @ref PadStrideInfo.
-     */
-    ARM_COMPUTE_DEPRECATED_REL(20.11)
-    void configure(const CLCompileContext &compile_context, const ICLTensor *input, const ICLTensor *weights, const ICLTensor *biases, ICLTensor *output, const PadStrideInfo &conv_info);
-    /** Static function to check if given info will lead to a valid configuration of @ref CLLocallyConnectedLayer
-     *
-     * @param[in] input     Input tensor info. 3 lower dimensions represent a single input [width, height, IFM],
-     *                      while every optional dimension from 4 and above represent a batch of inputs.
-     *                      Data types supported: F32.
-     * @param[in] weights   Weights tensor info. Weights are 5D tensor with dimensions [kernel_x, kernel_y, IFM, OFM, num_patches]. Data type supported:Same as @p input.
-     * @param[in] biases    Biases tensor info. Shared biases supported. Biases are 2D tensor with dimensions [OFM, num_patches]. Data type supported:Same as @p input.
-     * @param[in] output    Output tensor info. 3 lower dimensions represent a single output [width, height, OFM], while the rest represent batch of outputs.
-     *                      Data types supported: Same as @p input.
-     * @param[in] conv_info Contains padding and stride information described in @ref PadStrideInfo.
-     *
-     * @return a status
-     */
-    static Status validate(const ITensorInfo *input, const ITensorInfo *weights, const ITensorInfo *biases, const ITensorInfo *output, const PadStrideInfo &conv_info);
-
-    // Inherited methods overridden:
-    void run() override;
-    void prepare() override;
-
-private:
-    MemoryGroup                                             _memory_group;
-    std::unique_ptr<CLIm2ColKernel>                         _input_im2col_kernel;
-    std::unique_ptr<CLWeightsReshapeKernel>                 _weights_reshape_kernel;
-    std::unique_ptr<CLLocallyConnectedMatrixMultiplyKernel> _mm_kernel;
-    std::unique_ptr<CLCol2ImKernel>                         _output_col2im_kernel;
-    CLTensor                                                _input_im2col_reshaped;
-    CLTensor                                                _weights_reshaped;
-    CLTensor                                                _gemm_output;
-    bool                                                    _is_prepared;
-    const ICLTensor                                        *_original_weights;
-};
-}
-#endif /* ARM_COMPUTE_CLLOCALLYCONNECTEDLAYER_H */
diff --git a/arm_compute/runtime/NEON/NEFunctions.h b/arm_compute/runtime/NEON/NEFunctions.h
index 3952d499de..e25b03d915 100644
--- a/arm_compute/runtime/NEON/NEFunctions.h
+++ b/arm_compute/runtime/NEON/NEFunctions.h
@@ -103,7 +103,6 @@
 #include "arm_compute/runtime/NEON/functions/NELSTMLayerQuantized.h"
 #include "arm_compute/runtime/NEON/functions/NELaplacianPyramid.h"
 #include "arm_compute/runtime/NEON/functions/NELaplacianReconstruct.h"
-#include "arm_compute/runtime/NEON/functions/NELocallyConnectedLayer.h"
 #include "arm_compute/runtime/NEON/functions/NELogical.h"
 #include "arm_compute/runtime/NEON/functions/NEMagnitude.h"
 #include "arm_compute/runtime/NEON/functions/NEMaxUnpoolingLayer.h"
diff --git a/arm_compute/runtime/NEON/functions/NELocallyConnectedLayer.h b/arm_compute/runtime/NEON/functions/NELocallyConnectedLayer.h
deleted file mode 100644
index 86e6300130..0000000000
--- a/arm_compute/runtime/NEON/functions/NELocallyConnectedLayer.h
+++ /dev/null
@@ -1,111 +0,0 @@
-/*
- * Copyright (c) 2017-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_NELOCALLYCONNECTEDLAYER_H
-#define ARM_COMPUTE_NELOCALLYCONNECTEDLAYER_H
-
-#include "arm_compute/runtime/IFunction.h"
-
-#include "arm_compute/core/Types.h"
-#include "arm_compute/runtime/IMemoryManager.h"
-#include "arm_compute/runtime/MemoryGroup.h"
-#include "arm_compute/runtime/NEON/functions/NECol2Im.h"
-#include "arm_compute/runtime/NEON/functions/NEIm2Col.h"
-#include "arm_compute/runtime/Tensor.h"
-
-#include <memory>
-
-namespace arm_compute
-{
-class INETensor;
-class NEWeightsReshapeKernel;
-class NELocallyConnectedMatrixMultiplyKernel;
-
-/** Basic function to compute the locally connected layer. This function calls the following NEON kernels:
- *
- * -# @ref NEWeightsReshapeKernel (executed only once for each configuration)
- * -# @ref NEIm2ColKernel
- * -# @ref NELocallyConnectedMatrixMultiplyKernel
- * -# @ref NECol2ImKernel
- */
-class NELocallyConnectedLayer : public IFunction
-{
-public:
-    /** Default constructor */
-    NELocallyConnectedLayer(std::shared_ptr<IMemoryManager> memory_manager = nullptr);
-    /** Prevent instances of this class from being copied (As this class contains pointers) */
-    NELocallyConnectedLayer(const NELocallyConnectedLayer &) = delete;
-    /** Prevent instances of this class from being moved (As this class contains pointers) */
-    NELocallyConnectedLayer(NELocallyConnectedLayer &&) = delete;
-    /** Prevent instances of this class from being copied (As this class contains pointers) */
-    NELocallyConnectedLayer &operator=(const NELocallyConnectedLayer &) = delete;
-    /** Prevent instances of this class from being moved (As this class contains pointers) */
-    NELocallyConnectedLayer &operator=(NELocallyConnectedLayer &&) = delete;
-    /** Default destructor */
-    ~NELocallyConnectedLayer();
-    /** Set the input and output tensors.
-     *
-     * @param[in]  input     Source tensor. 3 lower dimensions represent a single input [width, height, IFM],
-     *                       while every optional dimension from 4 and above represent a batch of inputs.
-     *                       Data types supported: F16, F32.
-     * @param[in]  weights   Weights tensor. Weights are 5D tensor with dimensions [kernel_x, kernel_y, IFM, OFM, num_patches]. Data type supported:Same as @p input.
-     * @param[in]  biases    Biases tensor. Shared biases supported. Biases are 2D tensor with dimensions [OFM, num_patches]. Data type supported:Same as @p input.
-     * @param[out] output    Destination tensor. 3 lower dimensions represent a single output [width, height, OFM], while the rest represent batch of outputs.
-     *                       Data types supported: Same as @p input.
-     * @param[in]  conv_info Contains padding and stride information described in @ref PadStrideInfo.
-     */
-    ARM_COMPUTE_DEPRECATED_REL(20.11)
-    void configure(const ITensor *input, const ITensor *weights, const ITensor *biases, ITensor *output, const PadStrideInfo &conv_info);
-    /** Static function to check if given info will lead to a valid configuration of @ref NELocallyConnectedLayer
-     *
-     * @param[in] input     Input tensor info. 3 lower dimensions represent a single input [width, height, IFM],
-     *                      while every optional dimension from 4 and above represent a batch of inputs.
-     *                      Data types supported: F16, F32.
-     * @param[in] weights   Weights tensor info. Weights are 5D tensor with dimensions [kernel_x, kernel_y, IFM, OFM, num_patches]. Data type supported:Same as @p input.
-     * @param[in] biases    Biases tensor info. Shared biases supported. Biases are 2D tensor with dimensions [OFM, num_patches]. Data type supported:Same as @p input.
-     * @param[in] output    Output tensor info. 3 lower dimensions represent a single output [width, height, OFM], while the rest represent batch of outputs.
-     *                      Data types supported: Same as @p input.
-     * @param[in] conv_info Contains padding and stride information described in @ref PadStrideInfo.
-     *
-     * @return a status
-     */
-    static Status validate(const ITensorInfo *input, const ITensorInfo *weights, const ITensorInfo *biases, const ITensorInfo *output, const PadStrideInfo &conv_info);
-
-    // Inherited methods overridden:
-    void run() override;
-    void prepare() override;
-
-private:
-    MemoryGroup                                             _memory_group;
-    NEIm2Col                                                _input_im2col;
-    std::unique_ptr<NEWeightsReshapeKernel>                 _weights_reshape_kernel;
-    std::unique_ptr<NELocallyConnectedMatrixMultiplyKernel> _mm_kernel;
-    NECol2Im                                                _output_col2im;
-    Tensor                                                  _input_im2col_reshaped;
-    Tensor                                                  _weights_reshaped;
-    Tensor                                                  _gemm_output;
-    bool                                                    _is_prepared;
-    const ITensor                                          *_original_weights;
-};
-}
-#endif /* ARM_COMPUTE_NELOCALLYCONNECTEDLAYER_H */
diff --git a/docs/00_introduction.dox b/docs/00_introduction.dox
index bdfbdb89b9..94f73f8cdb 100644
--- a/docs/00_introduction.dox
+++ b/docs/00_introduction.dox
@@ -88,8 +88,11 @@ If there is more than one release in a month then an extra sequential number is
 
 v21.02 Public major release
  - Upgraded C++ standard to C++14
- - Removed kernels:
+ - Removed functions:
+   - NELocallyConnectedLayer / CLLocallyConnectedLayer
+ - Remove kernels:
    - NEGEMMMatrixVectorMultiplyKernel
+   - NELocallyConnectedMatrixMultiplyKernel / CLLocallyConnectedMatrixMultiplyKernel
 
 v20.11 Public major release
  - Various bug fixes.
@@ -130,7 +133,7 @@ v20.11 Public major release
    - @ref NEGEMMInterleave4x4Kernel
    - @ref NEDirectConvolutionLayerKernel
    - @ref NEScaleKernel
-   - @ref NELocallyConnectedMatrixMultiplyKernel
+   - NELocallyConnectedMatrixMultiplyKernel
    - @ref NEGEMMLowpOffsetContributionKernel
    - @ref NEGEMMTranspose1xWKernel
    - @ref NEPoolingLayerKernel
@@ -1200,7 +1203,7 @@ v17.06 Public major release
     - @ref CLBatchNormalizationLayerKernel / @ref CLBatchNormalizationLayer
     - @ref CLDepthConcatenateLayerKernel / CLDepthConcatenateLayer
     - @ref CLHOGOrientationBinningKernel @ref CLHOGBlockNormalizationKernel, @ref CLHOGDetectorKernel / @ref CLHOGDescriptor @ref CLHOGDetector @ref CLHOGGradient @ref CLHOGMultiDetection
-    - @ref CLLocallyConnectedMatrixMultiplyKernel / @ref CLLocallyConnectedLayer
+    - CLLocallyConnectedMatrixMultiplyKernel / CLLocallyConnectedLayer
     - @ref CLWeightsReshapeKernel / @ref CLConvolutionLayerReshapeWeights
  - New C++ kernels:
     - @ref CPPDetectionWindowNonMaximaSuppressionKernel
@@ -1208,7 +1211,7 @@ v17.06 Public major release
     - @ref NEBatchNormalizationLayerKernel / @ref NEBatchNormalizationLayer
     - @ref NEDepthConcatenateLayerKernel / NEDepthConcatenateLayer
     - @ref NEDirectConvolutionLayerKernel / @ref NEDirectConvolutionLayer
-    - @ref NELocallyConnectedMatrixMultiplyKernel / @ref NELocallyConnectedLayer
+    - NELocallyConnectedMatrixMultiplyKernel / NELocallyConnectedLayer
     - @ref NEWeightsReshapeKernel / @ref NEConvolutionLayerReshapeWeights
 
 v17.05 Public bug fixes release
diff --git a/src/core/CL/CLKernels.h b/src/core/CL/CLKernels.h
index b335372fa9..eea90eb599 100644
--- a/src/core/CL/CLKernels.h
+++ b/src/core/CL/CLKernels.h
@@ -103,7 +103,6 @@
 #include "src/core/CL/kernels/CLIntegralImageKernel.h"
 #include "src/core/CL/kernels/CLL2NormalizeLayerKernel.h"
 #include "src/core/CL/kernels/CLLKTrackerKernel.h"
-#include "src/core/CL/kernels/CLLocallyConnectedMatrixMultiplyKernel.h"
 #include "src/core/CL/kernels/CLMagnitudePhaseKernel.h"
 #include "src/core/CL/kernels/CLMaxUnpoolingLayerKernel.h"
 #include "src/core/CL/kernels/CLMeanStdDevKernel.h"
diff --git a/src/core/CL/cl_kernels/gemm.cl b/src/core/CL/cl_kernels/gemm.cl
index b6afb85aa4..6883aafee5 100644
--- a/src/core/CL/cl_kernels/gemm.cl
+++ b/src/core/CL/cl_kernels/gemm.cl
@@ -4379,72 +4379,4 @@ __kernel void gemm_ma_f16(TENSOR3D_DECLARATION(src),
     vstore8(out, 0, (__global half *)dst.ptr);
 }
 #endif // defined(ARM_COMPUTE_OPENCL_FP16_ENABLED)
-#endif // defined(BETA)
-
-#if defined(WIDTH_VECTOR_A)
-/** This OpenCL kernel computes the vector by matrix multiplication between each row of A (src0) and matrix B (src1) used for locally connected layer
- *
- * @note The width of A need to be passed at compile time using -DWIDTH_VECTOR_A
- *
- * @note The input A and matrix B must not be reshaped
- *
- * @param[in]  src0_ptr                           Pointer to the source matrix. Supported data types: F32
- * @param[in]  src0_stride_x                      Stride of the source matrix in X dimension (in bytes)
- * @param[in]  src0_step_x                        src_stride_x * number of elements along X processed per workitem(in bytes)
- * @param[in]  src0_stride_y                      Stride of the source matrix in Y dimension (in bytes)
- * @param[in]  src0_step_y                        src_stride_y * number of elements along Y processed per workitem(in bytes)
- * @param[in]  src0_offset_first_element_in_bytes The offset of the first element in the source matrix
- * @param[in]  src1_ptr                           Pointer to the source matrix. Supported data types: same as @p src0_ptr
- * @param[in]  src1_stride_x                      Stride of the source matrix in X dimension (in bytes)
- * @param[in]  src1_step_x                        src_stride_x * number of elements along X processed per workitem(in bytes)
- * @param[in]  src1_stride_y                      Stride of the source matrix in Y dimension (in bytes)
- * @param[in]  src1_step_y                        src_stride_y * number of elements along Y processed per workitem(in bytes)
- * @param[in]  src1_stride_z                      Stride of the source matrix in Z dimension (in bytes)
- * @param[in]  src1_step_z                        src_stride_z * number of elements along Z processed per workitem(in bytes)
- * @param[in]  src1_offset_first_element_in_bytes The offset of the first element in the source matrix
- * @param[out] dst_ptr                            Pointer to the destination matrix Supported data types: same as @p src0_ptr
- * @param[in]  dst_stride_x                       Stride of the destination matrix in X dimension (in bytes)
- * @param[in]  dst_step_x                         dst_gx_stride_x * number of elements along X processed per workitem(in bytes)
- * @param[in]  dst_stride_y                       Stride of the destination matrix in Y dimension (in bytes)
- * @param[in]  dst_step_y                         dst_gx_stride_y * number of elements along Y processed per workitem(in bytes)
- * @param[in]  dst_offset_first_element_in_bytes  The offset of the first element in the destination matrix
- */
-__kernel void gemm_lc_vm_f32(IMAGE_DECLARATION(src0),
-                             TENSOR3D_DECLARATION(src1),
-                             IMAGE_DECLARATION(dst))
-{
-    int idx = get_global_id(0) * 4;
-    int idy = get_global_id(1);
-
-    // Compute the address for the vector A and matrix B
-    int2 src_addr = ((int2)(src0_offset_first_element_in_bytes + src0_stride_y * idy, src1_offset_first_element_in_bytes + src1_stride_z * idy));
-    src_addr.s1 += idx * sizeof(float);
-
-    int end_row_vec_a = src_addr.s0 + (WIDTH_VECTOR_A * sizeof(float));
-
-    float4 acc = 0.0f;
-
-    for(; src_addr.s0 <= (end_row_vec_a - 2 * (int)sizeof(float)); src_addr += (int2)(2 * sizeof(float), 2 * src1_stride_y))
-    {
-        float2 a0 = vload2(0, (__global float *)(src0_ptr + src_addr.s0));
-        float4 b0 = vload4(0, (__global float *)(src1_ptr + src_addr.s1));
-        float4 b1 = vload4(0, (__global float *)(src1_ptr + src_addr.s1 + src1_stride_y));
-
-        acc += b0 * (float4)a0.s0;
-        acc += b1 * (float4)a0.s1;
-    }
-
-    for(; src_addr.s0 < end_row_vec_a; src_addr += (int2)(sizeof(float), src1_stride_y))
-    {
-        float  a0 = *((__global float *)(src0_ptr + src_addr.s0));
-        float4 b0 = vload4(0, (__global float *)(src1_ptr + src_addr.s1));
-
-        acc += b0 * (float4)a0;
-    }
-
-    // Compute destination address
-    Image dst = CONVERT_TO_IMAGE_STRUCT(dst);
-
-    vstore4(acc, 0, (__global float *)(offset(&dst, 0, 0)));
-}
-#endif // defined(WIDTH_VECTOR_A)
+#endif // defined(BETA)
\ No newline at end of file
diff --git a/src/core/CL/kernels/CLLocallyConnectedMatrixMultiplyKernel.cpp b/src/core/CL/kernels/CLLocallyConnectedMatrixMultiplyKernel.cpp
deleted file mode 100644
index 49e04c32c2..0000000000
--- a/src/core/CL/kernels/CLLocallyConnectedMatrixMultiplyKernel.cpp
+++ /dev/null
@@ -1,145 +0,0 @@
-/*
- * Copyright (c) 2017-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/CLLocallyConnectedMatrixMultiplyKernel.h"
-
-#include "arm_compute/core/CL/CLHelpers.h"
-#include "arm_compute/core/CL/ICLTensor.h"
-#include "arm_compute/core/Helpers.h"
-#include "arm_compute/core/Utils.h"
-#include "src/core/AccessWindowStatic.h"
-#include "src/core/CL/CLValidate.h"
-#include "src/core/helpers/WindowHelpers.h"
-
-namespace arm_compute
-{
-CLLocallyConnectedMatrixMultiplyKernel::CLLocallyConnectedMatrixMultiplyKernel()
-    : _input0(nullptr), _input1(nullptr), _output(nullptr)
-{
-}
-
-namespace
-{
-Status validate_arguments(const ITensorInfo *input0, const ITensorInfo *input1, const ITensorInfo *output)
-{
-    ARM_COMPUTE_RETURN_ERROR_ON_NULLPTR(input0, input1, output);
-    ARM_COMPUTE_RETURN_ERROR_ON_F16_UNSUPPORTED(input0);
-    ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input0, 1, DataType::F32);
-    ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_TYPES(input0, input1, output);
-    ARM_COMPUTE_RETURN_ERROR_ON(input0->dimension(0) != input1->dimension(1));
-
-    return Status{};
-}
-
-std::tuple<Status, Window> validate_and_configure_window(ITensorInfo *input0, ITensorInfo *input1, ITensorInfo *output)
-{
-    const unsigned int num_elems_processed_per_iteration_x = max_cl_vector_width / data_size_from_type(input0->data_type());
-
-    Window win = calculate_max_window(*output, Steps(num_elems_processed_per_iteration_x));
-
-    AccessWindowHorizontal input0_access(input0, 0, num_elems_processed_per_iteration_x);
-    AccessWindowHorizontal input1_access(input1, 0, num_elems_processed_per_iteration_x);
-    AccessWindowHorizontal output_access(output, 0, num_elems_processed_per_iteration_x);
-
-    bool window_changed = update_window_and_padding(win, input0_access, input1_access, output_access);
-
-    output_access.set_valid_region(win, ValidRegion(Coordinates(), output->tensor_shape()));
-
-    Status err = (window_changed) ? ARM_COMPUTE_CREATE_ERROR(ErrorCode::RUNTIME_ERROR, "Insufficient Padding!") : Status{};
-
-    return std::make_tuple(err, win);
-}
-} // namespace
-
-void CLLocallyConnectedMatrixMultiplyKernel::configure(const ICLTensor *input0, const ICLTensor *input1, ICLTensor *output)
-{
-    configure(CLKernelLibrary::get().get_compile_context(), input0, input1, output);
-}
-
-void CLLocallyConnectedMatrixMultiplyKernel::configure(const CLCompileContext &compile_context, const ICLTensor *input0, const ICLTensor *input1, ICLTensor *output)
-{
-    ARM_COMPUTE_ERROR_ON_NULLPTR(input0, input1, output);
-    ARM_COMPUTE_ERROR_THROW_ON(validate_arguments(input0->info(), input1->info(), output->info()));
-
-    _input0 = input0;
-    _input1 = input1;
-    _output = output;
-
-    cl::NDRange lws_hint;
-    if(output->info()->dimension(1) == 196)
-    {
-        lws_hint = cl::NDRange(1, 7);
-    }
-    else
-    {
-        lws_hint = cl::NDRange(8, 8);
-    }
-
-    std::ostringstream    mm_arguments;
-    std::set<std::string> build_opts;
-
-    mm_arguments << "-DWIDTH_VECTOR_A=" << input0->info()->dimension(0) << " ";
-    build_opts.emplace(mm_arguments.str());
-
-    // Create kernel
-    std::string data_type_name = lower_string(string_from_data_type(input0->info()->data_type()));
-    _kernel                    = create_kernel(compile_context, ("gemm_lc_vm_" + data_type_name), build_opts);
-
-    // Configure kernel window
-    auto win_config = validate_and_configure_window(input0->info(), input1->info(), output->info());
-
-    ARM_COMPUTE_ERROR_THROW_ON(std::get<0>(win_config));
-
-    ICLKernel::configure_internal(std::get<1>(win_config), lws_hint);
-}
-
-Status CLLocallyConnectedMatrixMultiplyKernel::validate(const ITensorInfo *input0, const ITensorInfo *input1, const ITensorInfo *output)
-{
-    ARM_COMPUTE_RETURN_ON_ERROR(validate_arguments(input0, input1, output));
-    ARM_COMPUTE_RETURN_ON_ERROR(std::get<0>(validate_and_configure_window(input0->clone().get(), input1->clone().get(), output->clone().get())));
-
-    return Status{};
-}
-
-void CLLocallyConnectedMatrixMultiplyKernel::run(const Window &window, cl::CommandQueue &queue)
-{
-    ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this);
-    ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW(ICLKernel::window(), window);
-
-    Window slice = window.first_slice_window_2D();
-
-    Window matrix_b_window;
-    matrix_b_window.use_tensor_dimensions(_input1->info()->tensor_shape());
-    Window slice_matrix_b = matrix_b_window.first_slice_window_3D();
-
-    do
-    {
-        unsigned int idx = 0;
-        add_2D_tensor_argument(idx, _input0, slice);
-        add_3D_tensor_argument(idx, _input1, slice_matrix_b);
-        add_2D_tensor_argument(idx, _output, slice);
-        enqueue(queue, *this, slice, lws_hint());
-    }
-    while(window.slide_window_slice_2D(slice));
-}
-} // namespace arm_compute
diff --git a/src/core/CL/kernels/CLLocallyConnectedMatrixMultiplyKernel.h b/src/core/CL/kernels/CLLocallyConnectedMatrixMultiplyKernel.h
deleted file mode 100644
index 5d0a22afa5..0000000000
--- a/src/core/CL/kernels/CLLocallyConnectedMatrixMultiplyKernel.h
+++ /dev/null
@@ -1,85 +0,0 @@
-/*
- * Copyright (c) 2017-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_CLLOCALLYCONNECTEDMATRIXMULTIPLYKERNEL_H
-#define ARM_COMPUTE_CLLOCALLYCONNECTEDMATRIXMULTIPLYKERNEL_H
-
-#include "src/core/CL/ICLKernel.h"
-
-namespace arm_compute
-{
-class ICLTensor;
-
-/** OpenCL kernel to multiply each row of first tensor with low 2 dimensions of second tensor.
- *
- * @attention The second input tensor must have at least 2 dimensions (matrix)
- *
- */
-class CLLocallyConnectedMatrixMultiplyKernel : public ICLKernel
-{
-public:
-    /** Default constructor */
-    CLLocallyConnectedMatrixMultiplyKernel();
-    /** Prevent instances of this class from being copied (As this class contains pointers) */
-    CLLocallyConnectedMatrixMultiplyKernel(const CLLocallyConnectedMatrixMultiplyKernel &) = delete;
-    /** Prevent instances of this class from being copied (As this class contains pointers) */
-    CLLocallyConnectedMatrixMultiplyKernel &operator=(const CLLocallyConnectedMatrixMultiplyKernel &) = delete;
-    /** Allow instances of this class to be moved */
-    CLLocallyConnectedMatrixMultiplyKernel(CLLocallyConnectedMatrixMultiplyKernel &&) = default;
-    /** Allow instances of this class to be moved */
-    CLLocallyConnectedMatrixMultiplyKernel &operator=(CLLocallyConnectedMatrixMultiplyKernel &&) = default;
-    /** Initialise the kernel's input, output and alpha
-     *
-     * @param[in]  input0 First input tensor. Data types supported: F32
-     * @param[in]  input1 Second input tensor. Data type supported: same as @p input0
-     * @param[out] output Output tensor to store the result. Data type supported: same as @p input0
-     */
-    void configure(const ICLTensor *input0, const ICLTensor *input1, ICLTensor *output);
-    /** Initialise the kernel's input, output and alpha
-     *
-     * @param[in]  compile_context The compile context to be used.
-     * @param[in]  input0          First input tensor. Data types supported: F32
-     * @param[in]  input1          Second input tensor. Data type supported: same as @p input0
-     * @param[out] output          Output tensor to store the result. Data type supported: same as @p input0
-     */
-    void configure(const CLCompileContext &compile_context, const ICLTensor *input0, const ICLTensor *input1, ICLTensor *output);
-    /** Static function to check if given info will lead to a valid configuration of @ref CLLocallyConnectedMatrixMultiplyKernel
-     *
-     * @param[in] input0 First input tensor info. Data types supported: F32
-     * @param[in] input1 Second input tensor info. Data type supported: same as @p input0
-     * @param[in] output Output tensor info. Data type supported: same as @p input0
-     *
-     * @return a status
-     */
-    static Status validate(const ITensorInfo *input0, const ITensorInfo *input1, const ITensorInfo *output);
-
-    // Inherited methods overridden:
-    void run(const Window &window, cl::CommandQueue &queue) override;
-
-private:
-    const ICLTensor *_input0;
-    const ICLTensor *_input1;
-    ICLTensor       *_output;
-};
-} // namespace arm_compute
-#endif /* ARM_COMPUTE_CLLOCALLYCONNECTEDMATRIXMULTIPLYKERNEL_H */
diff --git a/src/core/NEON/NEKernels.h b/src/core/NEON/NEKernels.h
index 88fb8d4023..091130c23a 100644
--- a/src/core/NEON/NEKernels.h
+++ b/src/core/NEON/NEKernels.h
@@ -99,7 +99,6 @@
 #include "src/core/NEON/kernels/NEIntegralImageKernel.h"
 #include "src/core/NEON/kernels/NEL2NormalizeLayerKernel.h"
 #include "src/core/NEON/kernels/NELKTrackerKernel.h"
-#include "src/core/NEON/kernels/NELocallyConnectedMatrixMultiplyKernel.h"
 #include "src/core/NEON/kernels/NELogicalKernel.h"
 #include "src/core/NEON/kernels/NEMagnitudePhaseKernel.h"
 #include "src/core/NEON/kernels/NEMaxUnpoolingLayerKernel.h"
diff --git a/src/core/NEON/kernels/NELocallyConnectedMatrixMultiplyKernel.cpp b/src/core/NEON/kernels/NELocallyConnectedMatrixMultiplyKernel.cpp
deleted file mode 100644
index f11694dee4..0000000000
--- a/src/core/NEON/kernels/NELocallyConnectedMatrixMultiplyKernel.cpp
+++ /dev/null
@@ -1,391 +0,0 @@
-/*
- * Copyright (c) 2017-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/NELocallyConnectedMatrixMultiplyKernel.h"
-
-#include "arm_compute/core/Error.h"
-#include "arm_compute/core/Helpers.h"
-#include "arm_compute/core/IAccessWindow.h"
-#include "arm_compute/core/ITensor.h"
-#include "arm_compute/core/TensorInfo.h"
-#include "arm_compute/core/Types.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/NEFixedPoint.h"
-#include "src/core/helpers/AutoConfiguration.h"
-#include "src/core/helpers/WindowHelpers.h"
-
-#include <arm_neon.h>
-#include <cstddef>
-#include <cstdint>
-#include <tuple>
-
-namespace arm_compute
-{
-class Coordinates;
-
-namespace
-{
-void vector_matrix_multiply_f16(const ITensor *input0, const ITensor *input1, ITensor *output, const Window &window, const ThreadInfo &info)
-{
-#ifdef __ARM_FEATURE_FP16_VECTOR_ARITHMETIC
-    const auto width_matrix_b  = static_cast<int>(output->info()->dimension(0));
-    const auto in_b_stride     = static_cast<int>(input1->info()->strides_in_bytes()[1] / data_size_from_type(input1->info()->data_type()));
-    const auto num_elems_vec_a = static_cast<int>(input0->info()->dimension(0));
-
-    // The implementation computes 16 elements per iteration
-    const int window_start_x = 16 * info.thread_id;
-    const int window_step_x  = 16 * info.num_threads;
-    // Make sure (window_end_x - window_start_x) is a multiple of window_step_x
-    const int window_end_x = ceil_to_multiple(width_matrix_b - window_start_x, window_step_x) + window_start_x;
-
-    Window win_out(window);
-    win_out.set(Window::DimX, Window::Dimension(window_start_x, window_end_x, window_step_x));
-
-    Window win_a(window);
-    win_a.set(Window::DimX, Window::Dimension(0, 1, 1));
-
-    Iterator ina(input0, win_a);
-    Iterator out(output, win_out);
-
-    execute_window_loop(win_out, [&](const Coordinates & id)
-    {
-        if(id.x() > width_matrix_b)
-        {
-            return;
-        }
-
-        float16x8_t acc0 = vdupq_n_f16(0.f);
-        float16x8_t acc1 = vdupq_n_f16(0.f);
-        float16x8_t acc2 = vdupq_n_f16(0.f);
-        float16x8_t acc3 = vdupq_n_f16(0.f);
-
-        auto vec_a    = reinterpret_cast<const float16_t *>(ina.ptr());
-        auto matrix_b = reinterpret_cast<const float16_t *>(input1->ptr_to_element(Coordinates(id[0], 0, id[1])));
-
-        const float16_t *vec_a_end_addr = vec_a + num_elems_vec_a;
-
-        for(; vec_a <= (vec_a_end_addr - 4);)
-        {
-            const float16x4_t a0l = vld1_f16(vec_a);
-
-            float16x8_t b00 = vld1q_f16(matrix_b);
-            float16x8_t b01 = vld1q_f16(matrix_b + 8 + 0 * in_b_stride);
-            float16x8_t b02 = vld1q_f16(matrix_b + 16 + 0 * in_b_stride);
-            float16x8_t b03 = vld1q_f16(matrix_b + 24 + 0 * in_b_stride);
-
-            float16x8_t b10 = vld1q_f16(matrix_b + 0 + 1 * in_b_stride);
-            float16x8_t b11 = vld1q_f16(matrix_b + 8 + 1 * in_b_stride);
-            float16x8_t b12 = vld1q_f16(matrix_b + 16 + 1 * in_b_stride);
-            float16x8_t b13 = vld1q_f16(matrix_b + 24 + 1 * in_b_stride);
-
-            acc0 = vaddq_f16(acc0, vmulq_lane_f16(b00, a0l, 0));
-            acc1 = vaddq_f16(acc1, vmulq_lane_f16(b01, a0l, 0));
-            acc2 = vaddq_f16(acc2, vmulq_lane_f16(b02, a0l, 0));
-            acc3 = vaddq_f16(acc3, vmulq_lane_f16(b03, a0l, 0));
-            acc0 = vaddq_f16(acc0, vmulq_lane_f16(b10, a0l, 1));
-            acc1 = vaddq_f16(acc1, vmulq_lane_f16(b11, a0l, 1));
-            acc2 = vaddq_f16(acc2, vmulq_lane_f16(b12, a0l, 1));
-            acc3 = vaddq_f16(acc3, vmulq_lane_f16(b13, a0l, 1));
-
-            matrix_b += 2 * in_b_stride;
-
-            b00 = vld1q_f16(matrix_b);
-            b01 = vld1q_f16(matrix_b + 8 + 0 * in_b_stride);
-            b02 = vld1q_f16(matrix_b + 16 + 0 * in_b_stride);
-            b03 = vld1q_f16(matrix_b + 24 + 0 * in_b_stride);
-            b10 = vld1q_f16(matrix_b + 0 + 1 * in_b_stride);
-            b11 = vld1q_f16(matrix_b + 8 + 1 * in_b_stride);
-            b12 = vld1q_f16(matrix_b + 16 + 1 * in_b_stride);
-            b13 = vld1q_f16(matrix_b + 24 + 1 * in_b_stride);
-
-            acc0 = vaddq_f16(acc0, vmulq_lane_f16(b00, a0l, 2));
-            acc1 = vaddq_f16(acc1, vmulq_lane_f16(b01, a0l, 2));
-            acc2 = vaddq_f16(acc2, vmulq_lane_f16(b02, a0l, 2));
-            acc3 = vaddq_f16(acc3, vmulq_lane_f16(b03, a0l, 2));
-            acc0 = vaddq_f16(acc0, vmulq_lane_f16(b10, a0l, 3));
-            acc1 = vaddq_f16(acc1, vmulq_lane_f16(b11, a0l, 3));
-            acc2 = vaddq_f16(acc2, vmulq_lane_f16(b12, a0l, 3));
-            acc3 = vaddq_f16(acc3, vmulq_lane_f16(b13, a0l, 3));
-
-            vec_a += 4;
-            matrix_b += 2 * in_b_stride;
-        }
-
-        for(; vec_a < vec_a_end_addr;)
-        {
-            const float16_t   a0  = *vec_a;
-            const float16x8_t b00 = vld1q_f16(matrix_b);
-            const float16x8_t b01 = vld1q_f16(matrix_b + 8 + 0 * in_b_stride);
-            const float16x8_t b02 = vld1q_f16(matrix_b + 16 + 0 * in_b_stride);
-            const float16x8_t b03 = vld1q_f16(matrix_b + 24 + 0 * in_b_stride);
-
-            acc0 = vaddq_f16(acc0, vmulq_n_f16(b00, a0));
-            acc1 = vaddq_f16(acc1, vmulq_n_f16(b01, a0));
-            acc2 = vaddq_f16(acc2, vmulq_n_f16(b02, a0));
-            acc3 = vaddq_f16(acc3, vmulq_n_f16(b03, a0));
-
-            vec_a += 1;
-            matrix_b += in_b_stride;
-        }
-
-        const auto vec_out = reinterpret_cast<float16_t *>(out.ptr());
-
-        vst1q_f16(vec_out + 0, acc0);
-        vst1q_f16(vec_out + 8, acc1);
-        vst1q_f16(vec_out + 16, acc2);
-        vst1q_f16(vec_out + 24, acc3);
-    },
-    ina, out);
-#else  /* __ARM_FEATURE_FP16_VECTOR_ARITHMETIC */
-    ARM_COMPUTE_UNUSED(input0);
-    ARM_COMPUTE_UNUSED(input1);
-    ARM_COMPUTE_UNUSED(output);
-    ARM_COMPUTE_UNUSED(window);
-    ARM_COMPUTE_UNUSED(info);
-    ARM_COMPUTE_ERROR("Not supported, recompile with -march=armv8.2-a+fp16+simd.");
-#endif /* __ARM_FEATURE_FP16_VECTOR_ARITHMETIC */
-}
-
-void vector_matrix_multiply_f32(const ITensor *input0, const ITensor *input1, ITensor *output, const Window &window, const ThreadInfo &info)
-{
-    const auto width_matrix_b  = static_cast<int>(output->info()->dimension(0));
-    const auto in_b_stride     = static_cast<int>(input1->info()->strides_in_bytes()[1] / data_size_from_type(input1->info()->data_type()));
-    const auto num_elems_vec_a = static_cast<int>(input0->info()->dimension(0));
-
-    // The implementation computes 16 elements per iteration
-    const int window_start_x = 16 * info.thread_id;
-    const int window_step_x  = 16 * info.num_threads;
-    // Make sure (window_end_x - window_start_x) is a multiple of window_step_x
-    const int window_end_x = ceil_to_multiple(width_matrix_b - window_start_x, window_step_x) + window_start_x;
-
-    Window win_out(window);
-    win_out.set(Window::DimX, Window::Dimension(window_start_x, window_end_x, window_step_x));
-
-    Window win_a(window);
-    win_a.set(Window::DimX, Window::Dimension(0, 0, 0));
-
-    Iterator ina(input0, win_a);
-    Iterator out(output, win_out);
-
-    execute_window_loop(win_out, [&](const Coordinates & id)
-    {
-        if(id.x() > width_matrix_b)
-        {
-            return;
-        }
-
-        float32x4_t acc0 = vdupq_n_f32(0.f);
-        float32x4_t acc1 = vdupq_n_f32(0.f);
-        float32x4_t acc2 = vdupq_n_f32(0.f);
-        float32x4_t acc3 = vdupq_n_f32(0.f);
-
-        auto vec_a    = reinterpret_cast<const float *>(ina.ptr());
-        auto matrix_b = reinterpret_cast<const float *>(input1->ptr_to_element(Coordinates(id[0], 0, id[1])));
-
-#if __arm__
-        asm volatile("PLD [%0, #128*4]" ::"r"(reinterpret_cast<const uint8_t *>(vec_a)));
-        asm volatile("PLD [%0, #128*4]" ::"r"(reinterpret_cast<const uint8_t *>(matrix_b)));
-        asm volatile("PLD [%0, #128*4]" ::"r"(reinterpret_cast<const uint8_t *>(matrix_b + in_b_stride)));
-#endif /* __arm__ */
-
-        const float *vec_a_end_addr = vec_a + num_elems_vec_a;
-
-        for(; vec_a <= (vec_a_end_addr - 4);)
-        {
-            float32x2_t a0l = vld1_f32(vec_a);
-
-            float32x4_t b00 = vld1q_f32(matrix_b + 0 + 0 * in_b_stride);
-            float32x4_t b01 = vld1q_f32(matrix_b + 4 + 0 * in_b_stride);
-            float32x4_t b02 = vld1q_f32(matrix_b + 8 + 0 * in_b_stride);
-            float32x4_t b03 = vld1q_f32(matrix_b + 12 + 0 * in_b_stride);
-
-            float32x4_t b10 = vld1q_f32(matrix_b + 0 + 1 * in_b_stride);
-            float32x4_t b11 = vld1q_f32(matrix_b + 4 + 1 * in_b_stride);
-            float32x4_t b12 = vld1q_f32(matrix_b + 8 + 1 * in_b_stride);
-            float32x4_t b13 = vld1q_f32(matrix_b + 12 + 1 * in_b_stride);
-
-#if __arm__
-            asm volatile("PLD [%0, #128*4]" ::"r"(reinterpret_cast<const uint8_t *>(vec_a)));
-            asm volatile("PLD [%0, #128*1]" ::"r"(reinterpret_cast<const uint8_t *>(matrix_b + 1 * in_b_stride)));
-            asm volatile("PLD [%0, #128*1]" ::"r"(reinterpret_cast<const uint8_t *>(matrix_b + 2 * in_b_stride)));
-            asm volatile("PLD [%0, #128*1]" ::"r"(reinterpret_cast<const uint8_t *>(matrix_b + 3 * in_b_stride)));
-            asm volatile("PLD [%0, #128*1]" ::"r"(reinterpret_cast<const uint8_t *>(matrix_b + 4 * in_b_stride)));
-#endif /* __arm__ */
-
-            acc0 = vmlaq_lane_f32(acc0, b00, a0l, 0);
-            acc1 = vmlaq_lane_f32(acc1, b01, a0l, 0);
-            acc2 = vmlaq_lane_f32(acc2, b02, a0l, 0);
-            acc3 = vmlaq_lane_f32(acc3, b03, a0l, 0);
-
-            acc0 = vmlaq_lane_f32(acc0, b10, a0l, 1);
-            acc1 = vmlaq_lane_f32(acc1, b11, a0l, 1);
-            acc2 = vmlaq_lane_f32(acc2, b12, a0l, 1);
-            acc3 = vmlaq_lane_f32(acc3, b13, a0l, 1);
-
-            vec_a += 2;
-            matrix_b += 2 * in_b_stride;
-
-            a0l = vld1_f32(vec_a);
-
-            b00 = vld1q_f32(matrix_b + 0 + 0 * in_b_stride);
-            b01 = vld1q_f32(matrix_b + 4 + 0 * in_b_stride);
-            b02 = vld1q_f32(matrix_b + 8 + 0 * in_b_stride);
-            b03 = vld1q_f32(matrix_b + 12 + 0 * in_b_stride);
-
-            b10 = vld1q_f32(matrix_b + 0 + 1 * in_b_stride);
-            b11 = vld1q_f32(matrix_b + 4 + 1 * in_b_stride);
-            b12 = vld1q_f32(matrix_b + 8 + 1 * in_b_stride);
-            b13 = vld1q_f32(matrix_b + 12 + 1 * in_b_stride);
-
-            acc0 = vmlaq_lane_f32(acc0, b00, a0l, 0);
-            acc1 = vmlaq_lane_f32(acc1, b01, a0l, 0);
-            acc2 = vmlaq_lane_f32(acc2, b02, a0l, 0);
-            acc3 = vmlaq_lane_f32(acc3, b03, a0l, 0);
-
-            acc0 = vmlaq_lane_f32(acc0, b10, a0l, 1);
-            acc1 = vmlaq_lane_f32(acc1, b11, a0l, 1);
-            acc2 = vmlaq_lane_f32(acc2, b12, a0l, 1);
-            acc3 = vmlaq_lane_f32(acc3, b13, a0l, 1);
-
-            vec_a += 2;
-            matrix_b += 2 * in_b_stride;
-        }
-
-        for(; vec_a < vec_a_end_addr;)
-        {
-            const float a0 = *vec_a;
-
-            const float32x4_t b00 = vld1q_f32(matrix_b + 0 + 0 * in_b_stride);
-            const float32x4_t b01 = vld1q_f32(matrix_b + 4 + 0 * in_b_stride);
-            const float32x4_t b02 = vld1q_f32(matrix_b + 8 + 0 * in_b_stride);
-            const float32x4_t b03 = vld1q_f32(matrix_b + 12 + 0 * in_b_stride);
-
-            acc0 = vmlaq_n_f32(acc0, b00, a0);
-            acc1 = vmlaq_n_f32(acc1, b01, a0);
-            acc2 = vmlaq_n_f32(acc2, b02, a0);
-            acc3 = vmlaq_n_f32(acc3, b03, a0);
-
-            vec_a += 1;
-            matrix_b += in_b_stride;
-        }
-
-        const auto vec_out = reinterpret_cast<float *>(out.ptr());
-
-        vst1q_f32(vec_out + 0, acc0);
-        vst1q_f32(vec_out + 4, acc1);
-        vst1q_f32(vec_out + 8, acc2);
-        vst1q_f32(vec_out + 12, acc3);
-    },
-    ina, out);
-}
-
-Status validate_arguments(const ITensorInfo *input0, const ITensorInfo *input1, const ITensorInfo *output)
-{
-    ARM_COMPUTE_RETURN_ERROR_ON_CPU_F16_UNSUPPORTED(input0);
-    ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input0, 1, DataType::F16, DataType::F32);
-    ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input1, 1, DataType::F16, DataType::F32);
-    ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(output, 1, DataType::F16, DataType::F32);
-    ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_TYPES(input0, input1, output);
-    ARM_COMPUTE_RETURN_ERROR_ON(input0->dimension(0) != input1->dimension(1));
-
-    return Status{};
-}
-
-std::tuple<Status, Window> validate_and_configure_window(ITensorInfo *input0, ITensorInfo *input1, ITensorInfo *output)
-{
-    constexpr unsigned int num_elems_processed_per_iteration_x = 16;
-
-    Window win = calculate_max_window(*output, Steps(num_elems_processed_per_iteration_x));
-
-    AccessWindowHorizontal input0_access(input0, 0, num_elems_processed_per_iteration_x);
-    AccessWindowHorizontal input1_access(input1, 0, num_elems_processed_per_iteration_x);
-    AccessWindowHorizontal output_access(output, 0, num_elems_processed_per_iteration_x);
-
-    bool window_changed = update_window_and_padding(win, input0_access, input1_access, output_access);
-
-    output_access.set_valid_region(win, ValidRegion(Coordinates(), output->tensor_shape()));
-
-    Status err = (window_changed) ? ARM_COMPUTE_CREATE_ERROR(ErrorCode::RUNTIME_ERROR, "Insufficient Padding!") : Status{};
-
-    return std::make_tuple(err, win);
-}
-} // namespace
-
-NELocallyConnectedMatrixMultiplyKernel::NELocallyConnectedMatrixMultiplyKernel()
-    : _input0(nullptr), _input1(nullptr), _output(nullptr)
-{
-}
-
-void NELocallyConnectedMatrixMultiplyKernel::configure(const ITensor *input0, const ITensor *input1, ITensor *output)
-{
-    ARM_COMPUTE_ERROR_ON_NULLPTR(input0, input1, output);
-    ARM_COMPUTE_ERROR_THROW_ON(validate_arguments(input0->info(), input1->info(), output->info()));
-
-    _input0 = input0;
-    _input1 = input1;
-    _output = output;
-
-    // Configure kernel window
-    auto win_config = validate_and_configure_window(input0->info(), input1->info(), output->info());
-
-    ARM_COMPUTE_ERROR_THROW_ON(std::get<0>(win_config));
-
-    INEKernel::configure(std::get<1>(win_config));
-}
-
-Status NELocallyConnectedMatrixMultiplyKernel::validate(const ITensorInfo *input0, const ITensorInfo *input1, const ITensorInfo *output)
-{
-    ARM_COMPUTE_RETURN_ON_ERROR(validate_arguments(input0, input1, output));
-    ARM_COMPUTE_RETURN_ON_ERROR(std::get<0>(validate_and_configure_window(input0->clone().get(), input1->clone().get(), output->clone().get())));
-
-    return Status{};
-}
-
-void NELocallyConnectedMatrixMultiplyKernel::run(const Window &window, const ThreadInfo &info)
-{
-    ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this);
-    ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW(INEKernel::window(), window);
-
-    switch(_input0->info()->data_type())
-    {
-        case DataType::F16:
-        {
-            vector_matrix_multiply_f16(_input0, _input1, _output, window, info);
-            break;
-        }
-        case DataType::F32:
-        {
-            vector_matrix_multiply_f32(_input0, _input1, _output, window, info);
-            break;
-        }
-        default:
-        {
-            ARM_COMPUTE_ERROR("Data type not supported");
-            break;
-        }
-    }
-}
-} // namespace arm_compute
diff --git a/src/core/NEON/kernels/NELocallyConnectedMatrixMultiplyKernel.h b/src/core/NEON/kernels/NELocallyConnectedMatrixMultiplyKernel.h
deleted file mode 100644
index 72093b4bb7..0000000000
--- a/src/core/NEON/kernels/NELocallyConnectedMatrixMultiplyKernel.h
+++ /dev/null
@@ -1,79 +0,0 @@
-/*
- * Copyright (c) 2017-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_NELOCALLYCONNECTEDMATRIXMULTIPLYKERNEL_H
-#define ARM_COMPUTE_NELOCALLYCONNECTEDMATRIXMULTIPLYKERNEL_H
-
-#include "src/core/NEON/INEKernel.h"
-
-namespace arm_compute
-{
-class ITensor;
-
-/** NEON kernel to multiply each row of first tensor with low 2 dimensions of second tensor. */
-class NELocallyConnectedMatrixMultiplyKernel : public INEKernel
-{
-public:
-    const char *name() const override
-    {
-        return "NELocallyConnectedMatrixMultiplyKernel";
-    }
-    /** Default constructor */
-    NELocallyConnectedMatrixMultiplyKernel();
-    /** Prevent instances of this class from being copied (As this class contains pointers) */
-    NELocallyConnectedMatrixMultiplyKernel(const NELocallyConnectedMatrixMultiplyKernel &) = delete;
-    /** Prevent instances of this class from being copied (As this class contains pointers) */
-    NELocallyConnectedMatrixMultiplyKernel &operator=(const NELocallyConnectedMatrixMultiplyKernel &) = delete;
-    /** Allow instances of this class to be moved */
-    NELocallyConnectedMatrixMultiplyKernel(NELocallyConnectedMatrixMultiplyKernel &&) = default;
-    /** Allow instances of this class to be moved */
-    NELocallyConnectedMatrixMultiplyKernel &operator=(NELocallyConnectedMatrixMultiplyKernel &&) = default;
-    /** Default destructor */
-    ~NELocallyConnectedMatrixMultiplyKernel() = default;
-    /** Initialise the kernel's input and output
-     *
-     * @param[in]  input0 First input tensor. Data types supported: F16, F32
-     * @param[in]  input1 Second input tensor containing the Matrix B. Data type supported: same as @p input0
-     * @param[out] output Output tensor to store the result of matrix multiplication. Data type supported: same as @p input0
-     */
-    void configure(const ITensor *input0, const ITensor *input1, ITensor *output);
-    /** Static function to check if given info will lead to a valid configuration of @ref NELocallyConnectedMatrixMultiplyKernel
-     *
-     * @param[in] input0 First input tensor info. Data types supported: F16, F32
-     * @param[in] input1 Second input tensor info. Data type supported: same as @p input0
-     * @param[in] output Output tensor info. Data type supported: same as @p input0
-     *
-     * @return a status
-     */
-    static Status validate(const ITensorInfo *input0, const ITensorInfo *input1, const ITensorInfo *output);
-
-    // Inherited methods overridden:
-    void run(const Window &window, const ThreadInfo &info) override;
-
-private:
-    const ITensor *_input0;
-    const ITensor *_input1;
-    ITensor       *_output;
-};
-} // namespace arm_compute
-#endif /* ARM_COMPUTE_NELOCALLYCONNECTEDMATRIXMULTIPLYKERNEL_H */
diff --git a/src/runtime/CL/functions/CLLocallyConnectedLayer.cpp b/src/runtime/CL/functions/CLLocallyConnectedLayer.cpp
deleted file mode 100644
index 3adae07095..0000000000
--- a/src/runtime/CL/functions/CLLocallyConnectedLayer.cpp
+++ /dev/null
@@ -1,226 +0,0 @@
-/*
- * Copyright (c) 2017-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 "arm_compute/runtime/CL/functions/CLLocallyConnectedLayer.h"
-
-#include "arm_compute/core/PixelValue.h"
-#include "arm_compute/core/Utils.h"
-#include "arm_compute/core/Validate.h"
-#include "arm_compute/runtime/CL/CLScheduler.h"
-#include "src/core/CL/kernels/CLCol2ImKernel.h"
-#include "src/core/CL/kernels/CLIm2ColKernel.h"
-#include "src/core/CL/kernels/CLLocallyConnectedMatrixMultiplyKernel.h"
-#include "src/core/CL/kernels/CLWeightsReshapeKernel.h"
-
-#include <cmath>
-#include <tuple>
-
-using namespace arm_compute;
-
-namespace
-{
-void calculate_shapes(const ITensorInfo *input, const ITensorInfo *weights, const ITensorInfo *biases, const ITensorInfo *output, const PadStrideInfo &conv_info,
-                      TensorShape &shape_wr, TensorShape &shape_im2col, TensorShape &shape_gemm)
-{
-    ARM_COMPUTE_UNUSED(output);
-
-    const unsigned int kernel_width  = weights->dimension(0);
-    const unsigned int kernel_height = weights->dimension(1);
-
-    bool has_bias = (biases != nullptr);
-
-    // Get convolved dimensions
-    unsigned int conv_w = 0;
-    unsigned int conv_h = 0;
-    std::tie(conv_w, conv_h) = scaled_dimensions(input->dimension(0),
-                                                 input->dimension(1),
-                                                 kernel_width,
-                                                 kernel_height,
-                                                 conv_info);
-
-    const size_t mat_weights_cols = weights->dimension(3);
-    const size_t mat_weights_rows = weights->dimension(0) * weights->dimension(1) * weights->dimension(2) + ((has_bias) ? 1 : 0);
-    const size_t mat_weights_num  = weights->dimension(4);
-
-    shape_wr = TensorShape(mat_weights_cols, mat_weights_rows, mat_weights_num);
-
-    const size_t mat_input_cols = mat_weights_rows;
-    const size_t mat_input_rows = conv_w * conv_h;
-
-    shape_im2col = input->tensor_shape();
-    if(shape_im2col.num_dimensions() >= 3)
-    {
-        shape_im2col.remove_dimension(2);
-    }
-    shape_im2col.set(0, mat_input_cols);
-    shape_im2col.set(1, mat_input_rows);
-
-    shape_gemm = shape_im2col;
-    shape_gemm.set(0, mat_weights_cols);
-    shape_gemm.set(1, mat_input_rows);
-}
-} // namespace
-
-CLLocallyConnectedLayer::CLLocallyConnectedLayer(std::shared_ptr<IMemoryManager> memory_manager)
-    : _memory_group(std::move(memory_manager)),
-      _input_im2col_kernel(std::make_unique<CLIm2ColKernel>()),
-      _weights_reshape_kernel(std::make_unique<CLWeightsReshapeKernel>()),
-      _mm_kernel(std::make_unique<CLLocallyConnectedMatrixMultiplyKernel>()),
-      _output_col2im_kernel(std::make_unique<CLCol2ImKernel>()),
-      _input_im2col_reshaped(),
-      _weights_reshaped(),
-      _gemm_output(),
-      _is_prepared(false),
-      _original_weights(nullptr)
-{
-}
-
-Status CLLocallyConnectedLayer::validate(const ITensorInfo *input, const ITensorInfo *weights, const ITensorInfo *biases, const ITensorInfo *output, const PadStrideInfo &conv_info)
-{
-    ARM_COMPUTE_RETURN_ERROR_ON_NULLPTR(input, weights, output);
-    ARM_COMPUTE_RETURN_ERROR_ON(weights->dimension(2) != input->dimension(2));
-    ARM_COMPUTE_RETURN_ERROR_ON(!conv_info.padding_is_symmetric());
-
-    bool has_bias = (biases != nullptr);
-
-    if(has_bias)
-    {
-        ARM_COMPUTE_RETURN_ERROR_ON(biases->dimension(0) != weights->dimension(3));
-        ARM_COMPUTE_RETURN_ERROR_ON(biases->num_dimensions() > 2);
-    }
-
-    const unsigned int kernel_width  = weights->dimension(0);
-    const unsigned int kernel_height = weights->dimension(1);
-
-    // Get convolved dimensions
-    unsigned int conv_w = 0;
-    unsigned int conv_h = 0;
-    std::tie(conv_w, conv_h) = scaled_dimensions(input->dimension(0), input->dimension(1), kernel_width, kernel_height,
-                                                 conv_info);
-
-    ARM_COMPUTE_RETURN_ERROR_ON_MSG((output->dimension(0) != conv_w) || (output->dimension(1) != conv_h), "Output shape does not match the expected one");
-    ARM_COMPUTE_RETURN_ERROR_ON_MSG(weights->dimension(4) != (conv_w * conv_h), "Weights shape does not match the expected one");
-
-    // Calculate intermediate buffer shapes
-    TensorShape shape_wr;
-    TensorShape shape_im2col;
-    TensorShape shape_gemm;
-    calculate_shapes(input, weights, biases, output, conv_info, shape_wr, shape_im2col, shape_gemm);
-
-    TensorInfo weights_reshaped_info(shape_wr, 1, weights->data_type());
-    TensorInfo input_im2col_reshaped_info(shape_im2col, 1, input->data_type());
-    TensorInfo gemm_output_info(shape_gemm, 1, input->data_type());
-
-    ARM_COMPUTE_RETURN_ON_ERROR(CLIm2ColKernel::validate(input, &input_im2col_reshaped_info, Size2D(kernel_width, kernel_height), conv_info, has_bias));
-    ARM_COMPUTE_RETURN_ON_ERROR(CLWeightsReshapeKernel::validate(weights, biases, &weights_reshaped_info));
-    ARM_COMPUTE_RETURN_ON_ERROR(CLLocallyConnectedMatrixMultiplyKernel::validate(&input_im2col_reshaped_info, &weights_reshaped_info, &gemm_output_info));
-    ARM_COMPUTE_RETURN_ON_ERROR(CLCol2ImKernel::validate(&gemm_output_info, output, Size2D(conv_w, conv_h)));
-
-    return Status{};
-}
-
-#pragma GCC diagnostic push
-#pragma GCC diagnostic ignored "-Wdeprecated-declarations"
-void CLLocallyConnectedLayer::configure(const ICLTensor *input, const ICLTensor *weights, const ICLTensor *biases, ICLTensor *output, const PadStrideInfo &conv_info)
-{
-    configure(CLKernelLibrary::get().get_compile_context(), input, weights, biases, output, conv_info);
-}
-#pragma GCC diagnostic pop
-
-void CLLocallyConnectedLayer::configure(const CLCompileContext &compile_context, const ICLTensor *input, const ICLTensor *weights, const ICLTensor *biases, ICLTensor *output,
-                                        const PadStrideInfo &conv_info)
-{
-    ARM_COMPUTE_ERROR_ON_NULLPTR(input, weights, output);
-    ARM_COMPUTE_ERROR_THROW_ON(CLLocallyConnectedLayer::validate(input->info(), weights->info(), biases == nullptr ? nullptr : biases->info(), output->info(), conv_info));
-
-    bool _has_bias    = (biases != nullptr);
-    _original_weights = weights;
-    _is_prepared      = false;
-
-    const unsigned int kernel_width  = weights->info()->dimension(0);
-    const unsigned int kernel_height = weights->info()->dimension(1);
-
-    // Get convolved dimensions
-    unsigned int conv_w = 0;
-    unsigned int conv_h = 0;
-    std::tie(conv_w, conv_h) = scaled_dimensions(input->info()->dimension(0), input->info()->dimension(1), kernel_width, kernel_height,
-                                                 conv_info);
-
-    // Calculate intermediate buffer shapes
-    TensorShape shape_wr;
-    TensorShape shape_im2col;
-    TensorShape shape_gemm;
-    calculate_shapes(input->info(), weights->info(), biases == nullptr ? nullptr : biases->info(), output->info(), conv_info, shape_wr, shape_im2col, shape_gemm);
-
-    _weights_reshaped.allocator()->init(TensorInfo(shape_wr, 1, weights->info()->data_type()));
-    _input_im2col_reshaped.allocator()->init(TensorInfo(shape_im2col, 1, input->info()->data_type()));
-    _gemm_output.allocator()->init(TensorInfo(shape_gemm, 1, input->info()->data_type()));
-
-    // Manage intermediate buffers
-    _memory_group.manage(&_input_im2col_reshaped);
-    _memory_group.manage(&_gemm_output);
-
-    // Configure kernels
-    _input_im2col_kernel->configure(compile_context, input, &_input_im2col_reshaped, Size2D(kernel_width, kernel_height), conv_info, _has_bias);
-    _weights_reshape_kernel->configure(compile_context, weights, biases, &_weights_reshaped);
-    _mm_kernel->configure(compile_context, &_input_im2col_reshaped, &_weights_reshaped, &_gemm_output);
-    _output_col2im_kernel->configure(compile_context, &_gemm_output, output, Size2D(conv_w, conv_h));
-
-    // Allocate intermediate tensors
-    _input_im2col_reshaped.allocator()->allocate();
-    _gemm_output.allocator()->allocate();
-
-    CLScheduler::get().tune_kernel_static(*_input_im2col_kernel);
-}
-
-void CLLocallyConnectedLayer::run()
-{
-    prepare();
-
-    MemoryGroupResourceScope scope_mg(_memory_group);
-
-    // Run input reshaping
-    CLScheduler::get().enqueue(*_input_im2col_kernel);
-
-    // Runs vector matrix multiply on reshaped matrices
-    CLScheduler::get().enqueue(*_mm_kernel);
-
-    // Reshape output matrix
-    CLScheduler::get().enqueue(*_output_col2im_kernel.get(), false);
-}
-
-void CLLocallyConnectedLayer::prepare()
-{
-    if(!_is_prepared)
-    {
-        ARM_COMPUTE_ERROR_ON(!_original_weights->is_used());
-
-        // Run weights reshaping and mark original weights tensor as unused
-        _weights_reshaped.allocator()->allocate();
-        CLScheduler::get().enqueue(*_weights_reshape_kernel);
-        _original_weights->mark_as_unused();
-
-        CLScheduler::get().queue().finish();
-        _is_prepared = true;
-    }
-}
diff --git a/src/runtime/NEON/functions/NELocallyConnectedLayer.cpp b/src/runtime/NEON/functions/NELocallyConnectedLayer.cpp
deleted file mode 100644
index c1164c3bee..0000000000
--- a/src/runtime/NEON/functions/NELocallyConnectedLayer.cpp
+++ /dev/null
@@ -1,203 +0,0 @@
-/*
- * Copyright (c) 2017-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 "arm_compute/runtime/NEON/functions/NELocallyConnectedLayer.h"
-
-#include "arm_compute/core/PixelValue.h"
-#include "arm_compute/core/Utils.h"
-#include "arm_compute/core/Validate.h"
-#include "arm_compute/runtime/NEON/NEScheduler.h"
-#include "src/core/NEON/kernels/NEIm2ColKernel.h"
-#include "src/core/NEON/kernels/NELocallyConnectedMatrixMultiplyKernel.h"
-#include "src/core/NEON/kernels/NEWeightsReshapeKernel.h"
-
-#include <cmath>
-#include <tuple>
-
-namespace arm_compute
-{
-namespace
-{
-void calculate_shapes(const ITensorInfo *input, const ITensorInfo *weights, const ITensorInfo *biases, const ITensorInfo *output, const PadStrideInfo &conv_info,
-                      TensorShape &shape_wr, TensorShape &shape_im2col, TensorShape &shape_gemm)
-{
-    ARM_COMPUTE_UNUSED(output);
-
-    const unsigned int kernel_width  = weights->dimension(0);
-    const unsigned int kernel_height = weights->dimension(1);
-
-    bool has_bias = (biases != nullptr);
-
-    // Get convolved dimensions
-    unsigned int conv_w = 0;
-    unsigned int conv_h = 0;
-    std::tie(conv_w, conv_h) = scaled_dimensions(input->dimension(0), input->dimension(1), kernel_width, kernel_height,
-                                                 conv_info);
-
-    const size_t mat_weights_cols = weights->dimension(3);
-    const size_t mat_weights_rows = weights->dimension(0) * weights->dimension(1) * weights->dimension(2) + ((has_bias) ? 1 : 0);
-    const size_t mat_weights_num  = weights->dimension(4);
-
-    shape_wr = TensorShape(mat_weights_cols, mat_weights_rows, mat_weights_num);
-
-    const size_t mat_input_cols = mat_weights_rows;
-    const size_t mat_input_rows = conv_w * conv_h;
-
-    shape_im2col = input->tensor_shape();
-    shape_im2col.set(0, mat_input_cols);
-    shape_im2col.set(1, mat_input_rows);
-    shape_im2col.set(2, 1);
-
-    shape_gemm = shape_im2col;
-    shape_gemm.set(0, mat_weights_cols);
-    shape_gemm.set(1, mat_input_rows);
-}
-} // namespace
-NELocallyConnectedLayer::~NELocallyConnectedLayer() = default;
-
-NELocallyConnectedLayer::NELocallyConnectedLayer(std::shared_ptr<IMemoryManager> memory_manager)
-    : _memory_group(std::move(memory_manager)), _input_im2col(), _weights_reshape_kernel(), _mm_kernel(), _output_col2im(), _input_im2col_reshaped(), _weights_reshaped(), _gemm_output(),
-      _is_prepared(false), _original_weights(nullptr)
-{
-}
-
-Status NELocallyConnectedLayer::validate(const ITensorInfo *input, const ITensorInfo *weights, const ITensorInfo *biases, const ITensorInfo *output, const PadStrideInfo &conv_info)
-{
-    ARM_COMPUTE_RETURN_ERROR_ON_NULLPTR(input, weights, output);
-    ARM_COMPUTE_RETURN_ERROR_ON(weights->dimension(2) != input->dimension(2));
-    ARM_COMPUTE_RETURN_ERROR_ON(!conv_info.padding_is_symmetric());
-
-    bool has_bias = (biases != nullptr);
-
-    if(has_bias)
-    {
-        ARM_COMPUTE_RETURN_ERROR_ON(biases->dimension(0) != weights->dimension(3));
-        ARM_COMPUTE_RETURN_ERROR_ON(biases->num_dimensions() > 2);
-    }
-
-    const unsigned int kernel_width  = weights->dimension(0);
-    const unsigned int kernel_height = weights->dimension(1);
-
-    // Get convolved dimensions
-    unsigned int conv_w = 0;
-    unsigned int conv_h = 0;
-    std::tie(conv_w, conv_h) = scaled_dimensions(input->dimension(0), input->dimension(1), kernel_width, kernel_height,
-                                                 conv_info);
-
-    ARM_COMPUTE_RETURN_ERROR_ON_MSG((output->dimension(0) != conv_w) || (output->dimension(1) != conv_h), "Output shape does not match the expected one");
-    ARM_COMPUTE_RETURN_ERROR_ON_MSG(weights->dimension(4) != (conv_w * conv_h), "Weights shape does not match the expected one");
-
-    // Calculate intermediate buffer shapes
-    TensorShape shape_wr;
-    TensorShape shape_im2col;
-    TensorShape shape_gemm;
-    calculate_shapes(input, weights, biases, output, conv_info, shape_wr, shape_im2col, shape_gemm);
-
-    TensorInfo weights_reshaped_info(shape_wr, 1, weights->data_type());
-    TensorInfo input_im2col_reshaped_info(shape_im2col, 1, input->data_type());
-    TensorInfo gemm_output_info(shape_gemm, 1, input->data_type());
-
-    ARM_COMPUTE_RETURN_ON_ERROR(NEIm2Col::validate(input, &input_im2col_reshaped_info, Size2D(kernel_width, kernel_height), conv_info, has_bias));
-    ARM_COMPUTE_RETURN_ON_ERROR(NEWeightsReshapeKernel::validate(weights, biases, &weights_reshaped_info));
-    ARM_COMPUTE_RETURN_ON_ERROR(NELocallyConnectedMatrixMultiplyKernel::validate(&input_im2col_reshaped_info, &weights_reshaped_info, &gemm_output_info));
-    ARM_COMPUTE_RETURN_ON_ERROR(NECol2Im::validate(&gemm_output_info, output, Size2D(conv_w, conv_h)));
-
-    return Status{};
-}
-
-void NELocallyConnectedLayer::configure(const ITensor *input, const ITensor *weights, const ITensor *biases, ITensor *output, const PadStrideInfo &conv_info)
-{
-    ARM_COMPUTE_ERROR_ON_NULLPTR(input, weights, output);
-    ARM_COMPUTE_ERROR_THROW_ON(NELocallyConnectedLayer::validate(input->info(), weights->info(), biases == nullptr ? nullptr : biases->info(), output->info(), conv_info));
-
-    bool _has_bias    = (biases != nullptr);
-    _is_prepared      = false;
-    _original_weights = weights;
-
-    const unsigned int kernel_width  = weights->info()->dimension(0);
-    const unsigned int kernel_height = weights->info()->dimension(1);
-
-    // Get convolved dimensions
-    unsigned int conv_w = 0;
-    unsigned int conv_h = 0;
-    std::tie(conv_w, conv_h) = scaled_dimensions(input->info()->dimension(0), input->info()->dimension(1), kernel_width, kernel_height,
-                                                 conv_info);
-
-    // Calculate intermediate buffer shapes
-    TensorShape shape_wr;
-    TensorShape shape_im2col;
-    TensorShape shape_gemm;
-    calculate_shapes(input->info(), weights->info(), biases == nullptr ? nullptr : biases->info(), output->info(), conv_info, shape_wr, shape_im2col, shape_gemm);
-
-    _weights_reshaped.allocator()->init(TensorInfo(shape_wr, 1, weights->info()->data_type()));
-    _input_im2col_reshaped.allocator()->init(TensorInfo(shape_im2col, 1, input->info()->data_type()));
-    _gemm_output.allocator()->init(TensorInfo(shape_gemm, 1, input->info()->data_type()));
-
-    // Manage intermediate buffers
-    _memory_group.manage(&_input_im2col_reshaped);
-    _memory_group.manage(&_gemm_output);
-
-    // Configure kernels
-    _input_im2col.configure(input, &_input_im2col_reshaped, Size2D(kernel_width, kernel_height), conv_info, _has_bias);
-    _weights_reshape_kernel = std::make_unique<NEWeightsReshapeKernel>();
-    _weights_reshape_kernel->configure(weights, biases, &_weights_reshaped);
-    _mm_kernel = std::make_unique<NELocallyConnectedMatrixMultiplyKernel>();
-    _mm_kernel->configure(&_input_im2col_reshaped, &_weights_reshaped, &_gemm_output);
-    _output_col2im.configure(&_gemm_output, output, Size2D(conv_w, conv_h));
-
-    // Allocate intermediate tensors
-    _input_im2col_reshaped.allocator()->allocate();
-    _gemm_output.allocator()->allocate();
-}
-
-void NELocallyConnectedLayer::run()
-{
-    prepare();
-
-    MemoryGroupResourceScope scope_mg(_memory_group);
-
-    // Run input reshaping
-    _input_im2col.run();
-
-    // Runs GEMM on reshaped matrices
-    NEScheduler::get().schedule(_mm_kernel.get(), Window::DimX);
-
-    // Reshape output matrix
-    _output_col2im.run();
-}
-
-void NELocallyConnectedLayer::prepare()
-{
-    if(!_is_prepared)
-    {
-        ARM_COMPUTE_ERROR_ON(!_original_weights->is_used());
-
-        // Run weights reshaping and mark original weights tensor as unused
-        _weights_reshaped.allocator()->allocate();
-        NEScheduler::get().schedule(_weights_reshape_kernel.get(), 3);
-        _original_weights->mark_as_unused();
-
-        _is_prepared = true;
-    }
-}
-} // namespace arm_compute