aboutsummaryrefslogtreecommitdiff
path: root/src/core/CL/cl_kernels/normalize_planar_yuv_layer.cl
diff options
context:
space:
mode:
authorMichele Di Giorgio <michele.digiorgio@arm.com>2018-09-12 10:18:54 +0100
committerAnthony Barbier <anthony.barbier@arm.com>2018-11-02 16:54:54 +0000
commitd63dfa2fc61a33b4e675ec6bc7458d8700174134 (patch)
tree2a85a8258aaa9a5762eb589f34b3f2868705dfb5 /src/core/CL/cl_kernels/normalize_planar_yuv_layer.cl
parent20c246a60869bada4051bd14eb9a3862be5330d7 (diff)
downloadComputeLibrary-d63dfa2fc61a33b4e675ec6bc7458d8700174134.tar.gz
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>
Diffstat (limited to 'src/core/CL/cl_kernels/normalize_planar_yuv_layer.cl')
-rw-r--r--src/core/CL/cl_kernels/normalize_planar_yuv_layer.cl12
1 files changed, 6 insertions, 6 deletions
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 dc6652449e..a105968a7b 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 @@ __kernel void normalize_planar_yuv_layer_nchw(TENSOR3D_DECLARATION(src),
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 @@ __kernel void normalize_planar_yuv_layer_nchw(TENSOR3D_DECLARATION(src),
(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 @@ __kernel void normalize_planar_yuv_layer_nhwc(TENSOR3D_DECLARATION(src),
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;