From efac7c63c045ea8fec17df54231f0d44611c7c78 Mon Sep 17 00:00:00 2001 From: Michele Di Giorgio Date: Wed, 16 May 2018 00:02:35 +0100 Subject: COMPMID-959: Fix CLChannelShuffle Change-Id: I1ea4db4e1ba37a736445ba991eeb08c247a6a61e Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/131393 Reviewed-by: Gian Marco Iodice Tested-by: Jenkins --- src/core/CL/cl_kernels/channel_shuffle.cl | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) (limited to 'src/core/CL') diff --git a/src/core/CL/cl_kernels/channel_shuffle.cl b/src/core/CL/cl_kernels/channel_shuffle.cl index d5cb10086e..26cee9ccdd 100644 --- a/src/core/CL/cl_kernels/channel_shuffle.cl +++ b/src/core/CL/cl_kernels/channel_shuffle.cl @@ -61,13 +61,13 @@ __kernel void channel_shuffle_nchw(TENSOR3D_DECLARATION(src), Tensor3D src = CONVERT_TO_TENSOR3D_STRUCT(src); Tensor3D dst = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(dst); - const uint curr_channel = get_global_id(2); // channel id of input - const uint group_id = curr_channel / NUM_GROUPS; // group id - const uint channel_id = curr_channel % NUM_GROUPS; // channel id within the group + const uint curr_channel = get_global_id(2); // channel id of input + const uint group_id = curr_channel / K; // group id + const uint channel_id = curr_channel % K; // channel id within the group const uint x = get_global_id(0) * BLOCK_SIZE; const uint y = get_global_id(1) * BLOCK_SIZE; - const uint z = channel_id * K + group_id; + const uint z = channel_id * NUM_GROUPS + group_id; // Load the NxN block TYPE u0 = VLOAD(BLOCK_SIZE)(0, (__global DATA_TYPE *)tensor3D_offset(&src, 0, 0, 0)); -- cgit v1.2.1