commit c878f1fdcc35c1d346fad6fbe975a55928f73187
author giuros01 <giuseppe.rossini@arm.com> Tue Jul 09 11:01:34 2019 +0100
committer Georgios Pinitas <georgios.pinitas@arm.com> Wed Jul 10 13:56:55 2019 +0000
tree d466a8d8c607a3fe1f6a0a2740ad4c892891a707
parent 01fd6cf41254995539d553baed6093f377871837

COMPMID-2457: Investigate DirectConvolution failures

We were not correctly handling the case where NumKernels > 1. In order
to handle this, I had to:
- Position the source pointer correctly (in the cl kernel)
- Position the weights pointer correctly (in the cl kernel)
- Set the correct num_elements_read_per_iteration_x and change the input access window
  accordingly

Change-Id: Ib2257b50930c822d3623f373dae04f188b46ee56
Signed-off-by: giuros01 <giuseppe.rossini@arm.com>
Reviewed-on: https://review.mlplatform.org/c/1498
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com>

src/core/CL/cl_kernels/direct_convolution9x9.cl

diff --git a/src/core/CL/cl_kernels/direct_convolution9x9.cl b/src/core/CL/cl_kernels/direct_convolution9x9.cl
index 22f437f..8d0417a 100644
--- a/src/core/CL/cl_kernels/direct_convolution9x9.cl
+++ b/src/core/CL/cl_kernels/direct_convolution9x9.cl
@@ -284,7 +284,9 @@
     const int id2 = get_global_id(2);
 
     __global uchar *weights_addr = (__global uchar *)tensor3D_offset(&weights, 0, 0, 0);
-    __global uchar *src_addr     = (__global uchar *)offset(&src, 0, 0) + ((id2 * STRIDE_Y) - PAD_TOP) * (int)src_stride_z;
+    __global uchar *src_addr     = (__global uchar *)offset(&src, 0, 0) - src_stride_x * id0 + ((id2 * STRIDE_Y) - PAD_TOP) * (int)src_stride_z;
+
+    weights_addr += id0 * weights_stride_w;
 
 #if(PAD_TOP == 1)
     const int coordy = id2 - PAD_TOP;