COMPMID-2077 Optimise CLL2NormalizeLayerKernel

Change-Id: If3866d90f84d62578b7e26f524fa4d7757a5970c
Signed-off-by: Usama Arif <usama.arif@arm.com>
Reviewed-on: https://review.mlplatform.org/c/906
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Michele Di Giorgio <michele.digiorgio@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Gian Marco Iodice <gianmarco.iodice@arm.com>

4 files changed, 26 insertions, 20 deletions

diff --git a/src/core/CL/cl_kernels/l2_normalize.cl b/src/core/CL/cl_kernels/l2_normalize.cl
index 5f66efbcc4..70b8b363fa 100644
--- a/src/core/CL/cl_kernels/l2_normalize.cl
+++ b/src/core/CL/cl_kernels/l2_normalize.cl
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2016-2018 ARM Limited.
+ * Copyright (c) 2016-2019 ARM Limited.
  *
  * SPDX-License-Identifier: MIT
  *
@@ -31,26 +31,32 @@
  * @param[in]  src_ptr                           Pointer to the source tensor. Supported data types: F16/F32
  * @param[in]  src_stride_x                      Stride of the source tensor in X dimension (in bytes)
  * @param[in]  src_step_x                        src_stride_x * number of elements along X processed per workitem(in bytes)
+ * @param[in]  src_stride_y                      Stride of the source tensor in Y dimension (in bytes)
+ * @param[in]  src_step_y                        src_stride_y * number of elements along X processed per workitem(in bytes)
  * @param[in]  src_offset_first_element_in_bytes The offset of the first element in the source tensor
  * @param[in]  sum_ptr                           Pointer to the source tensor. Supported data types: F16/F32
  * @param[in]  sum_stride_x                      Stride of the source tensor in X dimension (in bytes)
  * @param[in]  sum_step_x                        sum_stride_x * number of elements along X processed per workitem(in bytes)
+ * @param[in]  sum_stride_y                      Stride of the source tensor in Y dimension (in bytes)
+ * @param[in]  sum_step_y                        sum_stride_y * number of elements along Y processed per workitem(in bytes)
  * @param[in]  sum_offset_first_element_in_bytes The offset of the first element in the source tensor
  * @param[out] dst_ptr                           Pointer to the destination tensor. Supported data types: same as @p src_ptr
  * @param[in]  dst_stride_x                      Stride of the destination tensor in X dimension (in bytes)
  * @param[in]  dst_step_x                        dst_stride_x * number of elements along X processed per workitem(in bytes)
+ * @param[in]  dst_stride_y                      Stride of the destination tensor in Y dimension (in bytes)
+ * @param[in]  dst_step_y                        dst_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 tensor
  * @param[in]  epsilon                           Epsilon value
  */
 __kernel void l2_normalize_x(
-    VECTOR_DECLARATION(src),
-    VECTOR_DECLARATION(sum),
-    VECTOR_DECLARATION(dst),
+    IMAGE_DECLARATION(src),
+    IMAGE_DECLARATION(sum),
+    IMAGE_DECLARATION(dst),
     DATA_TYPE epsilon)
 {
-    Vector src = CONVERT_TO_VECTOR_STRUCT(src);
-    Vector sum = CONVERT_TO_VECTOR_STRUCT(sum);
-    Vector dst = CONVERT_TO_VECTOR_STRUCT(dst);
+    Image src = CONVERT_TO_IMAGE_STRUCT(src);
+    Image sum = CONVERT_TO_IMAGE_STRUCT(sum);
+    Image dst = CONVERT_TO_IMAGE_STRUCT(dst);
 
     VEC_DATA_TYPE(DATA_TYPE, 16)
     in = vload16(0, (__global DATA_TYPE *)src.ptr);
diff --git a/src/core/CL/kernels/CLL2NormalizeLayerKernel.cpp b/src/core/CL/kernels/CLL2NormalizeLayerKernel.cpp
index e33dab0f09..cb2e29449c 100644
--- a/src/core/CL/kernels/CLL2NormalizeLayerKernel.cpp
+++ b/src/core/CL/kernels/CLL2NormalizeLayerKernel.cpp
@@ -117,7 +117,7 @@ void CLL2NormalizeLayerKernel::configure(const ICLTensor *input, const ICLTensor
     {
         case 0:
             kernel_name = "x";
-            idx         = num_arguments_per_1D_tensor() * 3;
+            idx         = num_arguments_per_2D_tensor() * 3;
             break;
         case 1:
             kernel_name = "y";
@@ -169,17 +169,17 @@ void CLL2NormalizeLayerKernel::run(const Window &window, cl::CommandQueue &queue
         case 0:
         {
             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();
+            Window in_slice  = window.first_slice_window_2D();
+            Window sum_slice = window_sum.first_slice_window_2D();
             do
             {
                 unsigned int idx = 0;
-                add_1D_tensor_argument(idx, _input, in_slice);
-                add_1D_tensor_argument(idx, _sum, sum_slice);
-                add_1D_tensor_argument(idx, _output, in_slice);
+                add_2D_tensor_argument(idx, _input, in_slice);
+                add_2D_tensor_argument(idx, _sum, sum_slice);
+                add_2D_tensor_argument(idx, _output, in_slice);
                 enqueue(queue, *this, in_slice);
             }
-            while(window.slide_window_slice_1D(in_slice) && window.slide_window_slice_1D(sum_slice));
+            while(window.slide_window_slice_2D(in_slice) && window.slide_window_slice_2D(sum_slice));
         }
         break;
         case 1:
diff --git a/tests/benchmark/CL/L2NormalizeLayer.cpp b/tests/benchmark/CL/L2NormalizeLayer.cpp
index c88792c1ed..23bf612069 100644
--- a/tests/benchmark/CL/L2NormalizeLayer.cpp
+++ b/tests/benchmark/CL/L2NormalizeLayer.cpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2017 ARM Limited.
+ * Copyright (c) 2017-2019 ARM Limited.
  *
  * SPDX-License-Identifier: MIT
  *
@@ -49,10 +49,10 @@ using CLL2NormalizeLayerFixture = L2NormalizeLayerFixture<CLTensor, CLL2Normaliz
 TEST_SUITE(CL)
 
 REGISTER_FIXTURE_DATA_TEST_CASE(L2NormalizeLayer, CLL2NormalizeLayerFixture, framework::DatasetMode::ALL,
-                                framework::dataset::combine(framework::dataset::combine(datasets::SmallShapes(), data_types), framework::dataset::make("Axis", { 0 })));
+                                framework::dataset::combine(framework::dataset::combine(datasets::SmallShapes(), data_types), framework::dataset::make("Axis", { 0, 1, 2 })));
 TEST_SUITE(NIGHTLY)
 REGISTER_FIXTURE_DATA_TEST_CASE(L2NormalizeLayer, CLL2NormalizeLayerFixture, framework::DatasetMode::NIGHTLY,
-                                framework::dataset::combine(framework::dataset::combine(datasets::LargeShapes(), data_types), framework::dataset::make("Axis", { 0 })));
+                                framework::dataset::combine(framework::dataset::combine(datasets::LargeShapes(), data_types), framework::dataset::make("Axis", { 0, 1, 2 })));
 TEST_SUITE_END()
 TEST_SUITE_END()
 } // namespace benchmark
diff --git a/tests/benchmark/NEON/L2NormalizeLayer.cpp b/tests/benchmark/NEON/L2NormalizeLayer.cpp
index 14a9784312..164c64b2aa 100644
--- a/tests/benchmark/NEON/L2NormalizeLayer.cpp
+++ b/tests/benchmark/NEON/L2NormalizeLayer.cpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2017 ARM Limited.
+ * Copyright (c) 2017-2019 ARM Limited.
  *
  * SPDX-License-Identifier: MIT
  *
@@ -49,10 +49,10 @@ using NEL2NormalizeLayerFixture = L2NormalizeLayerFixture<Tensor, NEL2NormalizeL
 TEST_SUITE(NEON)
 
 REGISTER_FIXTURE_DATA_TEST_CASE(L2NormalizeLayer, NEL2NormalizeLayerFixture, framework::DatasetMode::ALL,
-                                framework::dataset::combine(framework::dataset::combine(datasets::SmallShapes(), data_types), framework::dataset::make("Axis", { 0 })));
+                                framework::dataset::combine(framework::dataset::combine(datasets::SmallShapes(), data_types), framework::dataset::make("Axis", { 0, 1, 2 })));
 TEST_SUITE(NIGHTLY)
 REGISTER_FIXTURE_DATA_TEST_CASE(L2NormalizeLayer, NEL2NormalizeLayerFixture, framework::DatasetMode::NIGHTLY,
-                                framework::dataset::combine(framework::dataset::combine(datasets::LargeShapes(), data_types), framework::dataset::make("Axis", { 0 })));
+                                framework::dataset::combine(framework::dataset::combine(datasets::LargeShapes(), data_types), framework::dataset::make("Axis", { 0, 1, 2 })));
 TEST_SUITE_END()
 TEST_SUITE_END()
 } // namespace benchmark