Update the heuristic for CLDepthwiseConvolutionNative kernel

- Use T_LOAD2D_INDIRECT macro instead of T_LOAD_NHWC_WITH_DILATION in
the depthwise convolution opencl kernels
- Update the heuristic for Arm® Mali™-G77

Resolves COMPMID-5716

Signed-off-by: Gian Marco Iodice <gianmarco.iodice@arm.com>
Change-Id: I32d375b220e04bf05f5d8f0af2231bde600f0665
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/8930
Benchmark: Arm Jenkins <bsgcomp@arm.com>
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Jakub Sujak <jakub.sujak@arm.com>
Reviewed-by: Viet-Hoa Do <viet-hoa.do@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>

3 files changed, 48 insertions, 6 deletions

diff --git a/src/core/CL/cl_kernels/nhwc/dwc_native_fp_nhwc.cl b/src/core/CL/cl_kernels/nhwc/dwc_native_fp_nhwc.cl
index dcbae220b6..6d64e270ef 100644
--- a/src/core/CL/cl_kernels/nhwc/dwc_native_fp_nhwc.cl
+++ b/src/core/CL/cl_kernels/nhwc/dwc_native_fp_nhwc.cl
@@ -108,7 +108,6 @@ __kernel void dwc_native_fp_nhwc(
 #define _IN0_A N0_A        // Cols tile A. It can be either 1 (for DEPTH_MULTIPLIER > 1) or N0 (for DEPTH_MULTIPLIER == 1)
 #define _IM0_B _IWEI_WIDTH // Rows tile B
 #define _IN0_B N0          // Cols tile B
-#define _IBOUNDARY_CHECK (!((WEI_WIDTH == 1 && WEI_HEIGHT == 1 && PAD_LEFT == 0 && PAD_TOP == 0 && M0 == 1)))
 
     const int cout = GET_SPATIAL_IDX(0, N0, PARTIAL_N0); // OFM
     const int xo   = GET_SPATIAL_IDX(1, M0, 0);          // WIDTH
@@ -146,8 +145,22 @@ __kernel void dwc_native_fp_nhwc(
             a[i].v = 0;
         })
 
+        TILE(int, 1, _IM0_A, mi);
+
+        LOOP_UNROLLING(int, xk_i, 0, 1, _IM0_A,
+        {
+            int x_s    = xi + xk_i * (DILATION_X);
+            int y_s    = yi + yk   * (DILATION_Y);
+            mi[0].s[xk_i] = x_s + y_s * SRC_WIDTH;
+            mi[0].s[xk_i] = mi[0].s[xk_i] + bout * (int)(SRC_WIDTH * SRC_HEIGHT);
+            mi[0].s[xk_i] = select(-1, mi[0].s[xk_i], x_s >= 0);
+            mi[0].s[xk_i] = select(-1, mi[0].s[xk_i], x_s < SRC_WIDTH);
+            mi[0].s[xk_i] = select(-1, mi[0].s[xk_i], y_s >= 0);
+            mi[0].s[xk_i] = select(-1, mi[0].s[xk_i], y_s < SRC_HEIGHT);
+        })
+
         // Load tile from the src tensor (TILE A)
-        T_LOAD_NHWC_WITH_DILATION(SRC_DATA_TYPE, 1, _IM0_A, _IN0_A, SRC_TENSOR_TYPE, src, bout, yi + yk * DILATION_Y, xi, (cout / DEPTH_MULTIPLIER), SRC_WIDTH, SRC_HEIGHT, DILATION_X, 1, _IBOUNDARY_CHECK, a);
+        T_LOAD2D_INDIRECT(SRC_DATA_TYPE, _IM0_A, _IN0_A, SRC_TENSOR_TYPE, src, (cout / DEPTH_MULTIPLIER), src_stride_y, mi, a);
 
         TILE(WEI_DATA_TYPE, _IM0_B, _IN0_B, b);
 
diff --git a/src/core/CL/cl_kernels/nhwc/dwc_native_quantized_nhwc.cl b/src/core/CL/cl_kernels/nhwc/dwc_native_quantized_nhwc.cl
index 2d255e5b61..e502d721d5 100644
--- a/src/core/CL/cl_kernels/nhwc/dwc_native_quantized_nhwc.cl
+++ b/src/core/CL/cl_kernels/nhwc/dwc_native_quantized_nhwc.cl
@@ -180,8 +180,22 @@ __kernel void dwc_native_quantized_nhwc(
         a[i].v = ZERO_VALUE;
     })
 
-    // Load tile from the src tensor (TILE A)
-    T_LOAD_NHWC_WITH_DILATION(SRC_DATA_TYPE, 1, _IM0_A, _IN0_A, SRC_TENSOR_TYPE, src, bout, yi + yk * DILATION_Y, xi, (cout / DEPTH_MULTIPLIER), src_w, src_h, DILATION_X, 1, _IBOUNDARY_CHECK, a);
+    TILE(int, 1, _IM0_A, my);
+
+    LOOP_UNROLLING(int, xk_i, 0, 1, _IM0_A,
+    {
+        int x_s    = xi + xk_i * (DILATION_X);
+        int y_s    = yi + yk   * (DILATION_Y);
+        my[0].s[xk_i] = x_s + y_s * SRC_WIDTH;
+        my[0].s[xk_i] = my[0].s[xk_i] + bout * (int)(SRC_WIDTH * SRC_HEIGHT);
+        my[0].s[xk_i] = select(-1, my[0].s[xk_i], x_s >= 0);
+        my[0].s[xk_i] = select(-1, my[0].s[xk_i], x_s < SRC_WIDTH);
+        my[0].s[xk_i] = select(-1, my[0].s[xk_i], y_s >= 0);
+        my[0].s[xk_i] = select(-1, my[0].s[xk_i], y_s < SRC_HEIGHT);
+    })
+
+    // Load tile from the src tensor
+    T_LOAD2D_INDIRECT(SRC_DATA_TYPE, _IM0_A, _IN0_A, SRC_TENSOR_TYPE, src, (cout / DEPTH_MULTIPLIER), src_stride_y, my, a);
 
     TILE(WEI_DATA_TYPE, _IM0_B, _IN0_B, b);
 
diff --git a/src/runtime/heuristics/dwc_native/ClDWCNativeDefaultConfigValhall.cpp b/src/runtime/heuristics/dwc_native/ClDWCNativeDefaultConfigValhall.cpp
index 49485c83a9..d0ade1bdd7 100644
--- a/src/runtime/heuristics/dwc_native/ClDWCNativeDefaultConfigValhall.cpp
+++ b/src/runtime/heuristics/dwc_native/ClDWCNativeDefaultConfigValhall.cpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2022 Arm Limited.
+ * Copyright (c) 2022-2023 Arm Limited.
  *
  * SPDX-License-Identifier: MIT
  *
@@ -283,15 +283,30 @@ DWCComputeKernelInfo ClDWCNativeDefaultConfigValhall::configure_G77_f16(const IT
         desc.n0 = adjust_vec_size(desc.n0, kernel_c);
 
         // Set m0 only if stride_x == 1 and dilation_x == 1
+        // m0 affects the number of rows to load from the input tensor. In fact, when depth_multiplier = 1, the number of rows
+        // loaded from the input tensors are -> kernel_width - (M0 - 1)
+        // The bigger the kernel_width, the smaller the M0 to avoid register spilling.
         if(conv_info.stride().first == 1 && dilation.x() == 1)
         {
+            // When the kernel width and kernel height are unit, it means that we have a pointwise multiplication. Therefore, M0 can be 1
             if((kernel_w >= 9) || (kernel_w == 1))
             {
                 desc.m0 = 1;
             }
             else
             {
-                desc.m0 = 2;
+                switch(kernel_w)
+                {
+                    case 3:
+                        desc.m0 = 4;
+                        break;
+                    case 5:
+                        desc.m0 = 3;
+                        break;
+                    default:
+                        desc.m0 = 2;
+                        break;
+                }
             }
         }
         else