COMPMID-: Implement L2NormalizationLayer.

Change-Id: I2223850b84826a87650729d287c57c22bcab7462
Reviewed-on: http://mpd-gerrit.cambridge.arm.com/82767
Tested-by: Kaizen <jeremy.johnson+kaizengerrit@arm.com>
Reviewed-by: Anthony Barbier <anthony.barbier@arm.com>

21 files changed, 1410 insertions, 0 deletions

diff --git a/arm_compute/core/NEON/NEKernels.h b/arm_compute/core/NEON/NEKernels.h
index a326f36b76..e49d04a39b 100644
--- a/arm_compute/core/NEON/NEKernels.h
+++ b/arm_compute/core/NEON/NEKernels.h
@@ -70,6 +70,7 @@
 #include "arm_compute/core/NEON/kernels/NEHistogramKernel.h"
 #include "arm_compute/core/NEON/kernels/NEIm2ColKernel.h"
 #include "arm_compute/core/NEON/kernels/NEIntegralImageKernel.h"
+#include "arm_compute/core/NEON/kernels/NEL2NormalizeKernel.h"
 #include "arm_compute/core/NEON/kernels/NELKTrackerKernel.h"
 #include "arm_compute/core/NEON/kernels/NELocallyConnectedMatrixMultiplyKernel.h"
 #include "arm_compute/core/NEON/kernels/NEMagnitudePhaseKernel.h"
@@ -82,6 +83,7 @@
 #include "arm_compute/core/NEON/kernels/NEPixelWiseMultiplicationKernel.h"
 #include "arm_compute/core/NEON/kernels/NEPoolingLayerKernel.h"
 #include "arm_compute/core/NEON/kernels/NEROIPoolingLayerKernel.h"
+#include "arm_compute/core/NEON/kernels/NEReductionOperationKernel.h"
 #include "arm_compute/core/NEON/kernels/NERemapKernel.h"
 #include "arm_compute/core/NEON/kernels/NEScaleKernel.h"
 #include "arm_compute/core/NEON/kernels/NEScharr3x3Kernel.h"
diff --git a/arm_compute/core/NEON/kernels/NEL2NormalizeKernel.h b/arm_compute/core/NEON/kernels/NEL2NormalizeKernel.h
new file mode 100644
index 0000000000..545aeaa1d3
--- /dev/null
+++ b/arm_compute/core/NEON/kernels/NEL2NormalizeKernel.h
@@ -0,0 +1,70 @@
+/*
+ * Copyright (c) 2017 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_NEL2NORMALIZEKERNEL_H__
+#define __ARM_COMPUTE_NEL2NORMALIZEKERNEL_H__
+
+#include "arm_compute/core/NEON/INEKernel.h"
+
+namespace arm_compute
+{
+class ITensor;
+
+/** Interface for performing a L2 normalize on a given axis given the square sum of it in this axis */
+class NEL2NormalizeKernel : public INEKernel
+{
+public:
+    /** Default constructor */
+    NEL2NormalizeKernel();
+    /** Prevent instances of this class from being copied (As this class contains pointers) */
+    NEL2NormalizeKernel(const NEL2NormalizeKernel &) = delete;
+    /** Prevent instances of this class from being copied (As this class contains pointers) */
+    NEL2NormalizeKernel &operator=(const NEL2NormalizeKernel &) = delete;
+    /** Allow instances of this class to be moved */
+    NEL2NormalizeKernel(NEL2NormalizeKernel &&) = default;
+    /** Allow instances of this class to be moved */
+    NEL2NormalizeKernel &operator=(NEL2NormalizeKernel &&) = default;
+    /** Default destructor */
+    ~NEL2NormalizeKernel() = default;
+    /** Set the input and output tensors.
+     *
+     * @param[in]  input   Source tensor. Data types supported: F32.
+     * @param[in]  sum     Sum values tensor. Data types supported: same as @p input.
+     * @param[out] output  Destination tensor. Data types supported: same as @p input.
+     * @param[in]  axis    Dimension along which to reduce. Supported reduction axis : 0
+     * @param[in]  epsilon Lower bound value for the normalization.
+     */
+    void configure(const ITensor *input, const ITensor *sum, ITensor *output, unsigned int axis, float epsilon);
+
+    // Inherited methods overridden:
+    void run(const Window &window) override;
+
+private:
+    const ITensor *_input;
+    const ITensor *_sum;
+    ITensor       *_output;
+    unsigned int   _axis;
+    float          _epsilon;
+};
+}
+#endif /*__ARM_COMPUTE_NEL2NORMALIZEKERNEL_H__ */
diff --git a/arm_compute/core/NEON/kernels/NEReductionOperationKernel.h b/arm_compute/core/NEON/kernels/NEReductionOperationKernel.h
new file mode 100644
index 0000000000..147e8351d6
--- /dev/null
+++ b/arm_compute/core/NEON/kernels/NEReductionOperationKernel.h
@@ -0,0 +1,69 @@
+/*
+ * Copyright (c) 2017 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_NEREDUCTIONOPERATIONKERNEL_H__
+#define __ARM_COMPUTE_NEREDUCTIONOPERATIONKERNEL_H__
+
+#include "arm_compute/core/NEON/INEKernel.h"
+
+namespace arm_compute
+{
+class ITensor;
+
+/** NEON kernel to perform a reduction operation */
+class NEReductionOperationKernel : public INEKernel
+{
+public:
+    /** Default constructor */
+    NEReductionOperationKernel();
+    /** Prevent instances of this class from being copied (As this class contains pointers) */
+    NEReductionOperationKernel(const NEReductionOperationKernel &) = delete;
+    /** Prevent instances of this class from being copied (As this class contains pointers) */
+    NEReductionOperationKernel &operator=(const NEReductionOperationKernel &) = delete;
+    /** Allow instances of this class to be moved */
+    NEReductionOperationKernel(NEReductionOperationKernel &&) = default;
+    /** Allow instances of this class to be moved */
+    NEReductionOperationKernel &operator=(NEReductionOperationKernel &&) = default;
+    /** Default destructor */
+    ~NEReductionOperationKernel() = default;
+    /** Set the source, destination of the kernel
+     *
+     * @param[in]  input  Source tensor. Data type supported: F32.
+     * @param[out] output Destination tensor.Data types supported: same as @p input.
+     * @param[in]  axis   Axis along which to reduce. Supported reduction axis : 0
+     * @param[in]  op     Reduction operation to perform.
+     */
+    void configure(const ITensor *input, ITensor *output, unsigned int axis, ReductionOperation op);
+    // Inherited methods overridden:
+    void run(const Window &window) override;
+    BorderSize border_size() const override;
+
+private:
+    const ITensor     *_input;
+    ITensor           *_output;
+    unsigned int       _reduction_axis;
+    ReductionOperation _op;
+    BorderSize         _border_size;
+};
+}
+#endif /*__ARM_COMPUTE_NEREDUCTIONOPERATIONKERNEL_H__ */
diff --git a/arm_compute/core/Types.h b/arm_compute/core/Types.h
index 2d3b3d6f66..3092976149 100644
--- a/arm_compute/core/Types.h
+++ b/arm_compute/core/Types.h
@@ -342,6 +342,12 @@ enum class NonLinearFilterFunction : unsigned
     MAX    = 2, /**< Non linear dilate. */
 };
 
+/** Available reduction operations */
+enum class ReductionOperation
+{
+    SUM_SQUARE, /**< Sum of squares */
+};
+
 /** The normalization type used for the normalization layer */
 enum class NormType
 {
diff --git a/arm_compute/runtime/NEON/NEFunctions.h b/arm_compute/runtime/NEON/NEFunctions.h
index e60dda1b53..42580ceefb 100644
--- a/arm_compute/runtime/NEON/NEFunctions.h
+++ b/arm_compute/runtime/NEON/NEFunctions.h
@@ -67,6 +67,7 @@
 #include "arm_compute/runtime/NEON/functions/NEHarrisCorners.h"
 #include "arm_compute/runtime/NEON/functions/NEHistogram.h"
 #include "arm_compute/runtime/NEON/functions/NEIntegralImage.h"
+#include "arm_compute/runtime/NEON/functions/NEL2Normalize.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"
@@ -82,6 +83,7 @@
 #include "arm_compute/runtime/NEON/functions/NEPixelWiseMultiplication.h"
 #include "arm_compute/runtime/NEON/functions/NEPoolingLayer.h"
 #include "arm_compute/runtime/NEON/functions/NEROIPoolingLayer.h"
+#include "arm_compute/runtime/NEON/functions/NEReductionOperation.h"
 #include "arm_compute/runtime/NEON/functions/NERemap.h"
 #include "arm_compute/runtime/NEON/functions/NEScale.h"
 #include "arm_compute/runtime/NEON/functions/NEScharr3x3.h"
diff --git a/arm_compute/runtime/NEON/functions/NEL2Normalize.h b/arm_compute/runtime/NEON/functions/NEL2Normalize.h
new file mode 100644
index 0000000000..1297b99e79
--- /dev/null
+++ b/arm_compute/runtime/NEON/functions/NEL2Normalize.h
@@ -0,0 +1,65 @@
+/*
+ * Copyright (c) 2017 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_NEL2NORMALIZE_H__
+#define __ARM_COMPUTE_NEL2NORMALIZE_H__
+
+#include "arm_compute/core/NEON/kernels/NEL2NormalizeKernel.h"
+#include "arm_compute/runtime/IFunction.h"
+#include "arm_compute/runtime/NEON/functions/NEReductionOperation.h"
+#include "arm_compute/runtime/Tensor.h"
+
+namespace arm_compute
+{
+class ITensor;
+
+/** Basic function to perform a L2 normalization on a given axis.
+ *
+ * This function runs the following kernels:
+ * -# @ref NEReductionOperation
+ * -# @ref NEL2NormalizeKernel
+ */
+class NEL2Normalize : public IFunction
+{
+public:
+    /** Constructor */
+    NEL2Normalize();
+    /** Set the input and output tensors.
+     *
+     * @param[in, out] input   Source tensor. Data types supported: F32. (Written to only for border_size != 0)
+     * @param[out]     output  Destination tensor. Data types supported: same as @p input.
+     * @param[in]      axis    Dimension along which to reduce. Supported reduction axis : 0
+     * @param[in]      epsilon Lower bound value for the normalization.
+     */
+    void configure(ITensor *input, ITensor *output, unsigned int axis, float epsilon = 1e-12);
+
+    // Inherited methods overridden:
+    void run() override;
+
+private:
+    NEReductionOperation _reduce_func;
+    NEL2NormalizeKernel  _normalize_kernel;
+    Tensor               _sumsq;
+};
+}
+#endif /* __ARM_COMPUTE_NEL2NORMALIZE_H__ */
diff --git a/arm_compute/runtime/NEON/functions/NEReductionOperation.h b/arm_compute/runtime/NEON/functions/NEReductionOperation.h
new file mode 100644
index 0000000000..9aa5228b16
--- /dev/null
+++ b/arm_compute/runtime/NEON/functions/NEReductionOperation.h
@@ -0,0 +1,66 @@
+/*
+ * Copyright (c) 2017 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_NEREDUCTIONOPERATION_H__
+#define __ARM_COMPUTE_NEREDUCTIONOPERATION_H__
+
+#include "arm_compute/runtime/IFunction.h"
+
+#include "arm_compute/core/NEON/kernels/NEFillBorderKernel.h"
+#include "arm_compute/core/NEON/kernels/NEReductionOperationKernel.h"
+#include "arm_compute/core/Types.h"
+
+namespace arm_compute
+{
+class ITensor;
+
+/** Basic function to simulate a reduction operation. This function calls the following NEON kernels:
+ *
+ * -# @ref NEFillBorderKernel
+ * -# @ref NEReductionOperationKernel
+ *
+ */
+class NEReductionOperation : public IFunction
+{
+public:
+    /** Default constructor */
+    NEReductionOperation();
+    /** Set the input and output tensors.
+     *
+     * @param[in, out] input  Source tensor. Data type supported: F32. (Written to only for border_size != 0)
+     * @param[out]     output Destination tensor. Data types supported: same as @p input.
+     * @param[in]      axis   Dimension along which to reduce. Supported reduction axis : 0
+     * @param[in]      op     Reduction operation to perform.
+     */
+    void configure(ITensor *input, ITensor *output, unsigned int axis, ReductionOperation op);
+
+    // Inherited methods overridden:
+    void run() override;
+
+private:
+    NEReductionOperationKernel _reduction_kernel;
+    NEFillBorderKernel         _fill_border_kernel;
+    size_t                     _window_split;
+};
+}
+#endif /* __ARM_COMPUTE_NEREDUCTIONOPERATION_H__ */
diff --git a/src/core/NEON/kernels/NEL2NormalizeKernel.cpp b/src/core/NEON/kernels/NEL2NormalizeKernel.cpp
new file mode 100644
index 0000000000..f3f58b6eb3
--- /dev/null
+++ b/src/core/NEON/kernels/NEL2NormalizeKernel.cpp
@@ -0,0 +1,125 @@
+/*
+ * Copyright (c) 2017 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/core/NEON/kernels/NEL2NormalizeKernel.h"
+
+#include "arm_compute/core/Error.h"
+#include "arm_compute/core/Helpers.h"
+#include "arm_compute/core/ITensor.h"
+#include "arm_compute/core/NEON/NEMath.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 <arm_neon.h>
+#include <cmath>
+
+using namespace arm_compute;
+
+namespace
+{
+void l2_normalize_X(const ITensor *in, const ITensor *sum, ITensor *out, float epsilon, const Window &window)
+{
+    Window window_sum(window);
+    window_sum.set(Window::DimX, Window::Dimension(0, 0, 0));
+
+    Window in_slice  = window.first_slice_window_1D();
+    Window sum_slice = window_sum.first_slice_window_1D();
+
+    do
+    {
+        Iterator input_it(in, in_slice);
+        Iterator sum_it(sum, sum_slice);
+        Iterator output_it(out, in_slice);
+
+        const float       sum_value           = *reinterpret_cast<const float *>(sum_it.ptr());
+        const float32x4_t vec_normalize_value = vdupq_n_f32(1.f / std::sqrt(std::max(sum_value, epsilon)));
+
+        execute_window_loop(in_slice, [&](const Coordinates & id)
+        {
+            const auto in_ptr  = reinterpret_cast<const float *>(input_it.ptr());
+            const auto out_ptr = reinterpret_cast<float *>(output_it.ptr());
+
+            vst1q_f32(out_ptr, vmulq_f32(vld1q_f32(in_ptr), vec_normalize_value));
+        },
+        input_it, output_it);
+    }
+    while(window.slide_window_slice_1D(in_slice) && window.slide_window_slice_1D(sum_slice));
+}
+} // namespace
+
+NEL2NormalizeKernel::NEL2NormalizeKernel()
+    : _input(nullptr), _sum(nullptr), _output(nullptr), _axis(0), _epsilon(1e-12)
+{
+}
+
+void NEL2NormalizeKernel::configure(const ITensor *input, const ITensor *sum, ITensor *output, unsigned int axis, float epsilon)
+{
+    ARM_COMPUTE_ERROR_ON_NULLPTR(input, sum, output);
+    ARM_COMPUTE_ERROR_ON_MSG(axis >= TensorShape::num_max_dimensions, "Normalization axis greater than max number of dimensions");
+    ARM_COMPUTE_ERROR_ON_MSG(axis > 0, "Unsupported normalization axis, Supported axis is 0");
+
+    // Output auto initialization if not yet initialized
+    auto_init_if_empty(*output->info(), input->info()->tensor_shape(), 1, input->info()->data_type(), input->info()->fixed_point_position());
+
+    ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::F32);
+    ARM_COMPUTE_ERROR_ON_MISMATCHING_DATA_TYPES(input, output, sum);
+    ARM_COMPUTE_ERROR_ON_MISMATCHING_SHAPES(input, output);
+
+    unsigned int num_elems_processed_per_iteration     = 16 / data_size_from_type(input->info()->data_type());
+    unsigned int num_elems_processed_per_iteration_sum = (axis == 0) ? 1 : num_elems_processed_per_iteration;
+
+    _input   = input;
+    _sum     = sum;
+    _output  = output;
+    _axis    = axis;
+    _epsilon = epsilon;
+
+    // Configure kernel window
+    Window                 win = calculate_max_window(*input->info(), Steps(num_elems_processed_per_iteration));
+    AccessWindowHorizontal input_access(input->info(), 0, num_elems_processed_per_iteration);
+    AccessWindowHorizontal sum_access(sum->info(), 0, num_elems_processed_per_iteration_sum);
+    AccessWindowHorizontal output_access(output->info(), 0, num_elems_processed_per_iteration);
+
+    update_window_and_padding(win, input_access, sum_access, output_access);
+
+    output_access.set_valid_region(win, input->info()->valid_region());
+
+    INEKernel::configure(win);
+}
+
+void NEL2NormalizeKernel::run(const Window &window)
+{
+    ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this);
+    ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW(INEKernel::window(), window);
+
+    switch(_axis)
+    {
+        case 0:
+            l2_normalize_X(_input, _sum, _output, _epsilon, window);
+            break;
+        default:
+            ARM_COMPUTE_ERROR("Unsupported normalization axis");
+    }
+}
diff --git a/src/core/NEON/kernels/NEReductionOperationKernel.cpp b/src/core/NEON/kernels/NEReductionOperationKernel.cpp
new file mode 100644
index 0000000000..5c8954609b
--- /dev/null
+++ b/src/core/NEON/kernels/NEReductionOperationKernel.cpp
@@ -0,0 +1,157 @@
+/*
+ * Copyright (c) 2017 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/core/NEON/kernels/NEReductionOperationKernel.h"
+
+#include "arm_compute/core/Coordinates.h"
+#include "arm_compute/core/Helpers.h"
+#include "arm_compute/core/IAccessWindow.h"
+#include "arm_compute/core/ITensor.h"
+#include "arm_compute/core/NEON/INEKernel.h"
+#include "arm_compute/core/NEON/NEMath.h"
+#include "arm_compute/core/Validate.h"
+
+#include <arm_neon.h>
+
+using namespace arm_compute;
+
+namespace
+{
+template <class F>
+class Reducer
+{
+public:
+    static void reduceX(const Window &window, const ITensor *input, ITensor *output, F f)
+    {
+        // Set out window
+        Window out_window(window);
+        out_window.set(Window::DimX, Window::Dimension(0, 0, 0));
+
+        // Get first input and output slices
+        Window in_slice  = window.first_slice_window_1D();
+        Window out_slice = out_window.first_slice_window_1D();
+
+        do
+        {
+            Iterator in(input, in_slice);
+            Iterator out(output, out_slice);
+
+            f(in, out, in_slice, out_slice);
+        }
+        while(window.slide_window_slice_1D(in_slice) && window.slide_window_slice_1D(out_slice));
+    }
+};
+
+struct SumsqOpX
+{
+    inline void operator()(Iterator &input, Iterator &output, Window &in_slice, Window &out_slice)
+    {
+        ARM_COMPUTE_UNUSED(out_slice);
+        float32x4_t vec_sum_value = vdupq_n_f32(0.f);
+
+        execute_window_loop(in_slice, [&](const Coordinates & id)
+        {
+            const auto        in_ptr       = reinterpret_cast<const float *>(input.ptr());
+            const float32x4_t vec_elements = vld1q_f32(in_ptr);
+            vec_sum_value                  = vaddq_f32(vmulq_f32(vec_elements, vec_elements), vec_sum_value);
+        },
+        input);
+
+        float32x2_t carry_addition = vpadd_f32(vget_high_f32(vec_sum_value), vget_low_f32(vec_sum_value));
+        carry_addition             = vpadd_f32(carry_addition, carry_addition);
+
+        *(reinterpret_cast<float *>(output.ptr())) = vget_lane_f32(carry_addition, 0);
+    }
+};
+
+void reduce_sumsq(const Window &window, const ITensor *input, ITensor *output, unsigned int axis)
+{
+    switch(axis)
+    {
+        case 0:
+            return Reducer<SumsqOpX>::reduceX(window, input, output, SumsqOpX());
+        default:
+            ARM_COMPUTE_ERROR("Unsupported reduction axis");
+    }
+}
+} // namespace
+
+NEReductionOperationKernel::NEReductionOperationKernel()
+    : _input(nullptr), _output(nullptr), _reduction_axis(0), _op(ReductionOperation::SUM_SQUARE), _border_size()
+{
+}
+
+BorderSize NEReductionOperationKernel::border_size() const
+{
+    return _border_size;
+}
+
+void NEReductionOperationKernel::configure(const ITensor *input, ITensor *output, unsigned int axis, ReductionOperation op)
+{
+    ARM_COMPUTE_ERROR_ON_NULLPTR(input, output);
+    ARM_COMPUTE_ERROR_ON_MSG(axis >= TensorShape::num_max_dimensions, "Reduction axis greater than max number of dimensions");
+    ARM_COMPUTE_ERROR_ON_MSG(axis > 0, "Unsupported reduction axis, Supported axis is 0");
+
+    // Calculate output shape and set if empty
+    TensorShape output_shape{ input->info()->tensor_shape() };
+    output_shape.set(axis, 1);
+
+    // Output auto initialization if not yet initialized
+    auto_init_if_empty(*output->info(), output_shape, 1, input->info()->data_type(), input->info()->fixed_point_position());
+
+    ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::F32);
+    ARM_COMPUTE_ERROR_ON_MISMATCHING_DATA_TYPES(input, output);
+    ARM_COMPUTE_ERROR_ON_MISMATCHING_DIMENSIONS(output->info()->tensor_shape(), output_shape);
+
+    unsigned int num_elems_processed_per_iteration = 16 / data_size_from_type(input->info()->data_type());
+
+    _input       = input;
+    _output      = output;
+    _border_size = (axis == 0) ? BorderSize(0, num_elems_processed_per_iteration - (input->info()->dimension(0) % num_elems_processed_per_iteration), 0, 0) : BorderSize();
+    _op          = op;
+
+    // Configure kernel window
+    Window                 win = calculate_max_window(*input->info(), Steps(num_elems_processed_per_iteration));
+    AccessWindowHorizontal input_access(input->info(), 0, num_elems_processed_per_iteration);
+    AccessWindowHorizontal output_access(output->info(), 0, num_elems_processed_per_iteration);
+
+    update_window_and_padding(win, input_access, output_access);
+    output_access.set_valid_region(win, ValidRegion(Coordinates(), output->info()->tensor_shape()));
+
+    INEKernel::configure(win);
+}
+
+void NEReductionOperationKernel::run(const Window &window)
+{
+    ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this);
+    ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW(INEKernel::window(), window);
+
+    switch(_op)
+    {
+        case ReductionOperation::SUM_SQUARE:
+            reduce_sumsq(window, _input, _output, _reduction_axis);
+            break;
+        default:
+            ARM_COMPUTE_ERROR("Unsupported reduction operation.");
+    }
+}
diff --git a/src/runtime/NEON/functions/NEL2Normalize.cpp b/src/runtime/NEON/functions/NEL2Normalize.cpp
new file mode 100644
index 0000000000..378d78e3f3
--- /dev/null
+++ b/src/runtime/NEON/functions/NEL2Normalize.cpp
@@ -0,0 +1,50 @@
+/*
+ * Copyright (c) 2017 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/NEL2Normalize.h"
+
+#include "arm_compute/core/Helpers.h"
+#include "arm_compute/runtime/NEON/NEScheduler.h"
+
+using namespace arm_compute;
+
+NEL2Normalize::NEL2Normalize()
+    : _reduce_func(), _normalize_kernel(), _sumsq()
+{
+}
+
+void NEL2Normalize::configure(ITensor *input, ITensor *output, unsigned int axis, float epsilon)
+{
+    // Configure Kernels
+    _reduce_func.configure(input, &_sumsq, axis, ReductionOperation::SUM_SQUARE);
+    _normalize_kernel.configure(input, &_sumsq, output, axis, epsilon);
+
+    // Allocate intermediate tensors
+    _sumsq.allocator()->allocate();
+}
+
+void NEL2Normalize::run()
+{
+    _reduce_func.run();
+    NEScheduler::get().schedule(&_normalize_kernel, Window::DimY);
+}
diff --git a/src/runtime/NEON/functions/NEReductionOperation.cpp b/src/runtime/NEON/functions/NEReductionOperation.cpp
new file mode 100644
index 0000000000..45c3e5dbc9
--- /dev/null
+++ b/src/runtime/NEON/functions/NEReductionOperation.cpp
@@ -0,0 +1,84 @@
+/*
+ * Copyright (c) 2017 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/NEReductionOperation.h"
+
+#include "arm_compute/core/Helpers.h"
+#include "arm_compute/runtime/NEON/NEScheduler.h"
+
+using namespace arm_compute;
+
+namespace
+{
+/** Define dimension to split the window
+ *
+ * @param[in] axis Reduction axis
+ *
+ * @return The dimension to split the window
+ */
+size_t reduction_window_split_dimension(unsigned int axis)
+{
+    switch(axis)
+    {
+        case 0:
+            return Window::DimY;
+        default:
+            ARM_COMPUTE_ERROR("Unsupported reduction axis");
+    }
+}
+BorderMode reduction_operation_border_mode(ReductionOperation op)
+{
+    switch(op)
+    {
+        case ReductionOperation::SUM_SQUARE:
+            return BorderMode::CONSTANT;
+        default:
+            return BorderMode::CONSTANT;
+    }
+}
+} // namespace
+
+NEReductionOperation::NEReductionOperation()
+    : _reduction_kernel(), _fill_border_kernel(), _window_split(0)
+{
+}
+
+void NEReductionOperation::configure(ITensor *input, ITensor *output, unsigned int axis, ReductionOperation op)
+{
+    ARM_COMPUTE_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::F32);
+
+    // Configure reduction kernel
+    _reduction_kernel.configure(input, output, axis, op);
+    _window_split = reduction_window_split_dimension(axis);
+
+    // Configure fill border kernel
+    BorderSize fill_border_size = (axis == 0) ? _reduction_kernel.border_size() : BorderSize();
+    BorderMode fill_border_mode = reduction_operation_border_mode(op);
+    _fill_border_kernel.configure(input, fill_border_size, fill_border_mode, PixelValue(0));
+}
+
+void NEReductionOperation::run()
+{
+    NEScheduler::get().schedule(&_fill_border_kernel, Window::DimY);
+    NEScheduler::get().schedule(&_reduction_kernel, _window_split);
+}
diff --git a/tests/TypePrinter.h b/tests/TypePrinter.h
index 5f7313552e..020b559c6f 100644
--- a/tests/TypePrinter.h
+++ b/tests/TypePrinter.h
@@ -198,6 +198,21 @@ inline ::std::ostream &operator<<(::std::ostream &os, const ConvertPolicy &polic
     return os;
 }
 
+/** Formatted output of the Reduction Operations. */
+inline ::std::ostream &operator<<(::std::ostream &os, const ReductionOperation &op)
+{
+    switch(op)
+    {
+        case ReductionOperation::SUM_SQUARE:
+            os << "SUM_SQUARE";
+            break;
+        default:
+            ARM_COMPUTE_ERROR("NOT_SUPPORTED!");
+    }
+
+    return os;
+}
+
 /** Formatted output of the activation function type. */
 inline ::std::ostream &operator<<(::std::ostream &os, const ActivationLayerInfo::ActivationFunction &act_function)
 {
@@ -240,6 +255,13 @@ inline ::std::ostream &operator<<(::std::ostream &os, const ActivationLayerInfo:
     return os;
 }
 
+inline std::string to_string(const ReductionOperation &op)
+{
+    std::stringstream str;
+    str << op;
+    return str.str();
+}
+
 inline std::string to_string(const ActivationLayerInfo::ActivationFunction &function)
 {
     std::stringstream str;
diff --git a/tests/datasets_new/ReductionOperationDataset.h b/tests/datasets_new/ReductionOperationDataset.h
new file mode 100644
index 0000000000..e46b8b1118
--- /dev/null
+++ b/tests/datasets_new/ReductionOperationDataset.h
@@ -0,0 +1,50 @@
+/*
+ * Copyright (c) 2017 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_TEST_REDUCTION_OPERATION_DATASET_H__
+#define __ARM_COMPUTE_TEST_REDUCTION_OPERATION_DATASET_H__
+
+#include "arm_compute/core/Types.h"
+#include "tests/TypePrinter.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace datasets
+{
+class ReductionOperations final : public framework::dataset::ContainerDataset<std::vector<ReductionOperation>>
+{
+public:
+    ReductionOperations()
+        : ContainerDataset("ReductionOperation",
+    {
+        ReductionOperation::SUM_SQUARE
+    })
+    {
+    }
+};
+} // namespace datasets
+} // namespace test
+} // namespace arm_compute
+#endif /* __ARM_COMPUTE_TEST_REDUCTION_OPERATION_DATASET_H__ */
diff --git a/tests/validation_new/CPP/L2Normalize.cpp b/tests/validation_new/CPP/L2Normalize.cpp
new file mode 100644
index 0000000000..f423702819
--- /dev/null
+++ b/tests/validation_new/CPP/L2Normalize.cpp
@@ -0,0 +1,88 @@
+/*
+ * Copyright (c) 2017 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 "L2Normalize.h"
+#include "ReductionOperation.h"
+
+#include "tests/validation_new/Helpers.h"
+
+#include <algorithm>
+#include <cmath>
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace reference
+{
+namespace
+{
+TensorShape get_output_shape(TensorShape shape, unsigned int axis)
+{
+    TensorShape output_shape(shape);
+    output_shape.set(axis, 1);
+    return output_shape;
+}
+} // namespace
+
+template <typename T>
+SimpleTensor<T> l2_normalize(const SimpleTensor<T> &src, unsigned int axis, float epsilon)
+{
+    // Create reference
+    SimpleTensor<T> dst{ src.shape(), src.data_type() };
+
+    // Reduce across given axis
+    SimpleTensor<T> sum = reduction_operation(src, get_output_shape(src.shape(), axis), axis, ReductionOperation::SUM_SQUARE);
+
+    // Compute reference
+    const int elems      = src.shape()[axis];
+    const int upper_dims = src.shape().total_size_upper(axis + 1);
+
+    for(int du = 0; du < upper_dims; ++du)
+    {
+        if(axis == 0)
+        {
+            const T *src_row_ptr         = src.data() + du * elems;
+            T       *dst_row_ptr         = dst.data() + du * elems;
+            const T  normalization_value = std::sqrt(std::max(sum[du], epsilon));
+            std::transform(src_row_ptr, src_row_ptr + elems, dst_row_ptr, [normalization_value](T val)
+            {
+                return val / normalization_value;
+            });
+        }
+        else
+        {
+            ARM_COMPUTE_ERROR("Unsupported normalization axis");
+        }
+    }
+
+    return dst;
+}
+
+template SimpleTensor<float> l2_normalize(const SimpleTensor<float> &src, unsigned int axis, float epsilon);
+} // namespace reference
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation_new/CPP/L2Normalize.h b/tests/validation_new/CPP/L2Normalize.h
new file mode 100644
index 0000000000..3e81b9a644
--- /dev/null
+++ b/tests/validation_new/CPP/L2Normalize.h
@@ -0,0 +1,44 @@
+/*
+ * Copyright (c) 2017 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_TEST_L2NORMALIZE_H__
+#define __ARM_COMPUTE_TEST_L2NORMALIZE_H__
+
+#include "tests/validation_new/Helpers.h"
+#include "tests/validation_new/SimpleTensor.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace reference
+{
+template <typename T>
+SimpleTensor<T> l2_normalize(const SimpleTensor<T> &src, unsigned int axis, float epsilon);
+} // namespace reference
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
+#endif /* __ARM_COMPUTE_TEST_L2NORMALIZE_H__ */
diff --git a/tests/validation_new/CPP/ReductionOperation.cpp b/tests/validation_new/CPP/ReductionOperation.cpp
new file mode 100644
index 0000000000..f4c5c21182
--- /dev/null
+++ b/tests/validation_new/CPP/ReductionOperation.cpp
@@ -0,0 +1,93 @@
+/*
+ * Copyright (c) 2017 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 "ReductionOperation.h"
+
+#include "tests/validation_new/Helpers.h"
+
+#include <algorithm>
+#include <cmath>
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace reference
+{
+namespace
+{
+template <typename T>
+struct square
+{
+    T operator()(const T &lhs, const T &rhs) const
+    {
+        return (lhs + rhs * rhs);
+    }
+};
+
+template <typename T>
+T reduce_operation(T *ptr, int reduce_elements, ReductionOperation op)
+{
+    switch(op)
+    {
+        case ReductionOperation::SUM_SQUARE:
+            return std::accumulate(ptr, ptr + reduce_elements, 0.f, square<T>());
+        default:
+            ARM_COMPUTE_ERROR("Unsupported reduction operation");
+    }
+}
+} // namespace
+
+template <typename T>
+SimpleTensor<T> reduction_operation(const SimpleTensor<T> &src, const TensorShape &dst_shape, unsigned int axis, ReductionOperation op)
+{
+    // Create reference
+    SimpleTensor<T> dst{ dst_shape, src.data_type() };
+
+    // Compute reference
+    const int reduce_elems = src.shape()[axis];
+    const int upper_dims   = src.shape().total_size_upper(axis + 1);
+
+    for(int du = 0; du < upper_dims; ++du)
+    {
+        if(axis == 0)
+        {
+            const T *src_row_ptr = src.data() + du * reduce_elems;
+            dst[du]              = reduce_operation(src_row_ptr, reduce_elems, op);
+        }
+        else
+        {
+            ARM_COMPUTE_ERROR("Unsupported reduction axis");
+        }
+    }
+
+    return dst;
+}
+
+template SimpleTensor<float> reduction_operation(const SimpleTensor<float> &src, const TensorShape &dst_shape, unsigned int axis, ReductionOperation op);
+} // namespace reference
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation_new/CPP/ReductionOperation.h b/tests/validation_new/CPP/ReductionOperation.h
new file mode 100644
index 0000000000..43a700d80a
--- /dev/null
+++ b/tests/validation_new/CPP/ReductionOperation.h
@@ -0,0 +1,44 @@
+/*
+ * Copyright (c) 2017 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_TEST_REDUCTION_OPERATION_H__
+#define __ARM_COMPUTE_TEST_REDUCTION_OPERATION_H__
+
+#include "tests/validation_new/Helpers.h"
+#include "tests/validation_new/SimpleTensor.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace reference
+{
+template <typename T>
+SimpleTensor<T> reduction_operation(const SimpleTensor<T> &src, const TensorShape &dst_shape, unsigned int axis, ReductionOperation op);
+} // namespace reference
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
+#endif /* __ARM_COMPUTE_TEST_FLOOR_H__ */
diff --git a/tests/validation_new/NEON/L2Normalize.cpp b/tests/validation_new/NEON/L2Normalize.cpp
new file mode 100644
index 0000000000..20804d2a94
--- /dev/null
+++ b/tests/validation_new/NEON/L2Normalize.cpp
@@ -0,0 +1,75 @@
+/*
+ * Copyright (c) 2017 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/core/Types.h"
+#include "arm_compute/runtime/NEON/functions/NEL2Normalize.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+#include "framework/Asserts.h"
+#include "framework/Macros.h"
+#include "framework/datasets/Datasets.h"
+#include "tests/NEON/Accessor.h"
+#include "tests/PaddingCalculator.h"
+#include "tests/datasets_new/ShapeDatasets.h"
+#include "tests/validation_new/Validation.h"
+#include "tests/validation_new/fixtures/L2NormalizeFixture.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace
+{
+/** Tolerance for float operations */
+constexpr RelativeTolerance tolerance_f32(0.00001f);
+} // namespace
+
+TEST_SUITE(NEON)
+TEST_SUITE(L2Normalize)
+
+template <typename T>
+using NEL2NormalizeFixture = L2NormalizeValidationFixture<Tensor, Accessor, NEL2Normalize, T>;
+
+TEST_SUITE(FP32)
+FIXTURE_DATA_TEST_CASE(RunSmall, NEL2NormalizeFixture<float>, framework::DatasetMode::PRECOMMIT,
+                       combine(combine(combine(datasets::SmallShapes(), framework::dataset::make("DataType", DataType::F32)), framework::dataset::make("Axis", { 0 })), framework::dataset::make("Epsilon", { 1e-12 })))
+{
+    // Validate output
+    validate(Accessor(_target), _reference, tolerance_f32);
+}
+
+FIXTURE_DATA_TEST_CASE(RunLarge, NEL2NormalizeFixture<float>, framework::DatasetMode::NIGHTLY,
+                       combine(combine(combine(datasets::LargeShapes(), framework::dataset::make("DataType", DataType::F32)), framework::dataset::make("Axis", { 0 })), framework::dataset::make("Epsilon", { 1e-12 })))
+{
+    // Validate output
+    validate(Accessor(_target), _reference, tolerance_f32);
+}
+TEST_SUITE_END()
+
+TEST_SUITE_END()
+TEST_SUITE_END()
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation_new/NEON/ReductionOperation.cpp b/tests/validation_new/NEON/ReductionOperation.cpp
new file mode 100644
index 0000000000..4048af09e7
--- /dev/null
+++ b/tests/validation_new/NEON/ReductionOperation.cpp
@@ -0,0 +1,75 @@
+/*
+ * Copyright (c) 2017 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/core/Types.h"
+#include "arm_compute/runtime/NEON/functions/NEReductionOperation.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "arm_compute/runtime/TensorAllocator.h"
+#include "framework/Asserts.h"
+#include "framework/Macros.h"
+#include "framework/datasets/Datasets.h"
+#include "tests/NEON/Accessor.h"
+#include "tests/PaddingCalculator.h"
+#include "tests/datasets_new/ReductionOperationDataset.h"
+#include "tests/datasets_new/ShapeDatasets.h"
+#include "tests/validation_new/Validation.h"
+#include "tests/validation_new/fixtures/ReductionOperationFixture.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+namespace
+{
+/** Tolerance for float operations */
+constexpr RelativeTolerance tolerance_f32(0.00001f);
+} // namespace
+
+TEST_SUITE(NEON)
+TEST_SUITE(ReductionOperation)
+
+template <typename T>
+using NEReductionOperationFixture = ReductionOperationValidationFixture<Tensor, Accessor, NEReductionOperation, T>;
+
+TEST_SUITE(FP32)
+FIXTURE_DATA_TEST_CASE(RunSmall, NEReductionOperationFixture<float>, framework::DatasetMode::PRECOMMIT,
+                       combine(combine(combine(datasets::SmallShapes(), framework::dataset::make("DataType", DataType::F32)), framework::dataset::make("Axis", { 0 })), datasets::ReductionOperations()))
+{
+    // Validate output
+    validate(Accessor(_target), _reference, tolerance_f32);
+}
+FIXTURE_DATA_TEST_CASE(RunLarge, NEReductionOperationFixture<float>, framework::DatasetMode::NIGHTLY,
+                       combine(combine(combine(datasets::LargeShapes(), framework::dataset::make("DataType", DataType::F32)), framework::dataset::make("Axis", { 0 })), datasets::ReductionOperations()))
+{
+    // Validate output
+    validate(Accessor(_target), _reference, tolerance_f32);
+}
+TEST_SUITE_END()
+
+TEST_SUITE_END()
+TEST_SUITE_END()
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
diff --git a/tests/validation_new/fixtures/L2NormalizeFixture.h b/tests/validation_new/fixtures/L2NormalizeFixture.h
new file mode 100644
index 0000000000..5102b93f51
--- /dev/null
+++ b/tests/validation_new/fixtures/L2NormalizeFixture.h
@@ -0,0 +1,107 @@
+/*
+ * Copyright (c) 2017 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_TEST_L2NORMALIZE_FIXTURE
+#define ARM_COMPUTE_TEST_L2NORMALIZE_FIXTURE
+
+#include "arm_compute/core/TensorShape.h"
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "framework/Asserts.h"
+#include "framework/Fixture.h"
+#include "tests/AssetsLibrary.h"
+#include "tests/Globals.h"
+#include "tests/IAccessor.h"
+#include "tests/validation_new/CPP/L2Normalize.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+template <typename TensorType, typename AccessorType, typename FunctionType, typename T>
+class L2NormalizeValidationFixture : public framework::Fixture
+{
+public:
+    template <typename...>
+    void setup(TensorShape shape, DataType data_type, unsigned int axis, float epsilon)
+    {
+        _target    = compute_target(shape, data_type, axis, epsilon);
+        _reference = compute_reference(shape, data_type, axis, epsilon);
+    }
+
+protected:
+    template <typename U>
+    void fill(U &&tensor)
+    {
+        library->fill_tensor_uniform(tensor, 0);
+    }
+
+    TensorType compute_target(const TensorShape &shape, DataType data_type, unsigned int axis, float epsilon)
+    {
+        // Create tensors
+        TensorType src = create_tensor<TensorType>(shape, data_type);
+        TensorType dst = create_tensor<TensorType>(shape, data_type);
+
+        // Create and configure function
+        FunctionType l2_norm_func;
+        l2_norm_func.configure(&src, &dst, axis, epsilon);
+
+        ARM_COMPUTE_EXPECT(src.info()->is_resizable(), framework::LogLevel::ERRORS);
+        ARM_COMPUTE_EXPECT(dst.info()->is_resizable(), framework::LogLevel::ERRORS);
+
+        // Allocate tensors
+        src.allocator()->allocate();
+        dst.allocator()->allocate();
+
+        ARM_COMPUTE_EXPECT(!src.info()->is_resizable(), framework::LogLevel::ERRORS);
+        ARM_COMPUTE_EXPECT(!dst.info()->is_resizable(), framework::LogLevel::ERRORS);
+
+        // Fill tensors
+        fill(AccessorType(src));
+
+        // Compute function
+        l2_norm_func.run();
+
+        return dst;
+    }
+
+    SimpleTensor<T> compute_reference(const TensorShape &shape, DataType data_type, unsigned int axis, float epsilon)
+    {
+        // Create reference
+        SimpleTensor<T> src{ shape, data_type };
+
+        // Fill reference
+        fill(src);
+
+        return reference::l2_normalize<T>(src, axis, epsilon);
+    }
+
+    TensorType      _target{};
+    SimpleTensor<T> _reference{};
+};
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
+#endif /* ARM_COMPUTE_TEST_L2NORMALIZE_FIXTURE */
diff --git a/tests/validation_new/fixtures/ReductionOperationFixture.h b/tests/validation_new/fixtures/ReductionOperationFixture.h
new file mode 100644
index 0000000000..416aadf8d7
--- /dev/null
+++ b/tests/validation_new/fixtures/ReductionOperationFixture.h
@@ -0,0 +1,116 @@
+/*
+ * Copyright (c) 2017 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_TEST_REDUCTION_OPERATION_FIXTURE
+#define ARM_COMPUTE_TEST_REDUCTION_OPERATION_FIXTURE
+
+#include "arm_compute/core/TensorShape.h"
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/Tensor.h"
+#include "framework/Asserts.h"
+#include "framework/Fixture.h"
+#include "tests/AssetsLibrary.h"
+#include "tests/Globals.h"
+#include "tests/IAccessor.h"
+#include "tests/validation_new/CPP/ReductionOperation.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace validation
+{
+template <typename TensorType, typename AccessorType, typename FunctionType, typename T>
+class ReductionOperationValidationFixture : public framework::Fixture
+{
+public:
+    template <typename...>
+    void setup(TensorShape shape, DataType data_type, unsigned int axis, ReductionOperation op)
+    {
+        const TensorShape output_shape = get_output_shape(shape, axis);
+        _target                        = compute_target(shape, output_shape, data_type, axis, op);
+        _reference                     = compute_reference(shape, output_shape, data_type, axis, op);
+    }
+
+protected:
+    template <typename U>
+    void fill(U &&tensor)
+    {
+        library->fill_tensor_uniform(tensor, 0);
+    }
+
+    TensorType compute_target(const TensorShape &src_shape, const TensorShape &dst_shape, DataType data_type, unsigned int axis, ReductionOperation op)
+    {
+        // Create tensors
+        TensorType src = create_tensor<TensorType>(src_shape, data_type);
+        TensorType dst = create_tensor<TensorType>(dst_shape, data_type);
+
+        // Create and configure function
+        FunctionType reduction_func;
+        reduction_func.configure(&src, &dst, axis, op);
+
+        ARM_COMPUTE_EXPECT(src.info()->is_resizable(), framework::LogLevel::ERRORS);
+        ARM_COMPUTE_EXPECT(dst.info()->is_resizable(), framework::LogLevel::ERRORS);
+
+        // Allocate tensors
+        src.allocator()->allocate();
+        dst.allocator()->allocate();
+
+        ARM_COMPUTE_EXPECT(!src.info()->is_resizable(), framework::LogLevel::ERRORS);
+        ARM_COMPUTE_EXPECT(!dst.info()->is_resizable(), framework::LogLevel::ERRORS);
+
+        // Fill tensors
+        fill(AccessorType(src));
+
+        // Compute function
+        reduction_func.run();
+
+        return dst;
+    }
+
+    SimpleTensor<T> compute_reference(const TensorShape &src_shape, const TensorShape &dst_shape, DataType data_type, unsigned int axis, ReductionOperation op)
+    {
+        // Create reference
+        SimpleTensor<T> src{ src_shape, data_type };
+
+        // Fill reference
+        fill(src);
+
+        return reference::reduction_operation<T>(src, dst_shape, axis, op);
+    }
+
+    TensorType      _target{};
+    SimpleTensor<T> _reference{};
+
+private:
+    TensorShape get_output_shape(TensorShape shape, unsigned int axis)
+    {
+        TensorShape output_shape(shape);
+        output_shape.set(axis, 1);
+        return output_shape;
+    }
+};
+} // namespace validation
+} // namespace test
+} // namespace arm_compute
+#endif /* ARM_COMPUTE_TEST_REDUCTION_OPERATION_FIXTURE */