commit d63dfa2fc61a33b4e675ec6bc7458d8700174134
author Michele Di Giorgio <michele.digiorgio@arm.com> Wed Sep 12 10:18:54 2018 +0100
committer Anthony Barbier <anthony.barbier@arm.com> Fri Nov 02 16:54:54 2018 +0000
tree 2a85a8258aaa9a5762eb589f34b3f2868705dfb5
parent 20c246a60869bada4051bd14eb9a3862be5330d7

COMPMID-1568: Add support for QASYMM8 to CLNormalizePlanarYUV

Change-Id: Id7ea6e7f57179478e5ba0e9231274e98fa089590
Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/148028
Tested-by: bsgcomp <bsgcomp@arm.com>
Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com>

src/core/CL/cl_kernels/normalize_planar_yuv_layer.cl

diff --git a/src/core/CL/cl_kernels/normalize_planar_yuv_layer.cl b/src/core/CL/cl_kernels/normalize_planar_yuv_layer.cl
index dc66524..a105968 100644
--- a/src/core/CL/cl_kernels/normalize_planar_yuv_layer.cl
+++ b/src/core/CL/cl_kernels/normalize_planar_yuv_layer.cl
@@ -27,7 +27,7 @@
 
 #define TYPE VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)
 
-/** Apply normalize_planar_yuv layer on tensors with NCHW format.
+/** Apply normalize_planar_yuv layer on tensors with NCHW data layout.
  *
  * @note Data type should be given as a preprocessor argument using -DDATA_TYPE=type. e.g. -DDATA_TYPE=float
  * @note Vector size should be given as a preprocessor argument using -DVEC_SIZE e.g. -DVEC_SIZE=8
@@ -70,8 +70,8 @@
 
     const uint current_slice = get_global_id(2) % NUM_CHANNELS;
 
-    const DATA_TYPE curr_mean = *((__global DATA_TYPE *)(mean.ptr + current_slice * mean.stride_x));
-    const DATA_TYPE curr_std  = *((__global DATA_TYPE *)(std.ptr + current_slice * std.stride_x));
+    const DATA_TYPE curr_mean = *((__global DATA_TYPE *)(mean.ptr + current_slice * sizeof(DATA_TYPE)));
+    const DATA_TYPE curr_std  = *((__global DATA_TYPE *)(std.ptr + current_slice * sizeof(DATA_TYPE)));
 
     TYPE data = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)src.ptr);
     TYPE res  = (data - curr_mean) / curr_std;
@@ -80,7 +80,7 @@
     (res, 0, (__global DATA_TYPE *)dst.ptr);
 }
 
-/** Apply normalize_planar_yuv layer on tensors with NHWC format.
+/** Apply normalize_planar_yuv layer on tensors with NHWC data layout.
  *
  * @note Data type should be given as a preprocessor argument using -DDATA_TYPE=type. e.g. -DDATA_TYPE=float
  * @note Vector size should be given as a preprocessor argument using -DVEC_SIZE e.g. -DVEC_SIZE=8
@@ -122,8 +122,8 @@
 
     const uint current_slice = get_global_id(0);
 
-    const TYPE curr_mean = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)(mean.ptr + current_slice * VEC_SIZE * mean.stride_x));
-    const TYPE curr_std  = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)(std.ptr + current_slice * VEC_SIZE * std.stride_x));
+    const TYPE curr_mean = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)(mean.ptr + current_slice * VEC_SIZE * sizeof(DATA_TYPE)));
+    const TYPE curr_std  = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)(std.ptr + current_slice * VEC_SIZE * sizeof(DATA_TYPE)));
 
     TYPE data = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)src.ptr);
     TYPE res  = (data - curr_mean) / curr_std;