diff options
| -rw-r--r-- | src/core/CL/cl_kernels/channel_shuffle.cl | 8 | ||||
| -rw-r--r-- | tests/validation/reference/ChannelShuffle.cpp | 10 |
2 files changed, 9 insertions, 9 deletions
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)); diff --git a/tests/validation/reference/ChannelShuffle.cpp b/tests/validation/reference/ChannelShuffle.cpp index 790eff53a0..c4d8d50e3d 100644 --- a/tests/validation/reference/ChannelShuffle.cpp +++ b/tests/validation/reference/ChannelShuffle.cpp @@ -57,13 +57,13 @@ SimpleTensor<T> channel_shuffle(const SimpleTensor<T> &src, int num_groups) { // Gather the group g block (of size channels_in_group * MxN) from output channels // g + 0 * G, g + 1 * G, g + 2 * G, g + G * (K - 1) etc. - const T *src_ptr = src_ref + g * MxN + n * num_channels * MxN; - T *dst_ptr = dst_ref + g * channels_in_group * MxN + n * num_channels * MxN; + const T *src_ptr = src_ref + g * channels_in_group * MxN + n * num_channels * MxN; + T *dst_ptr = dst_ref + g * MxN + n * num_channels * MxN; for(int i = 0; i < channels_in_group; ++i) { - std::copy(src_ptr + i * num_groups * MxN, - src_ptr + (i * num_groups + 1) * MxN, - dst_ptr + i * MxN); + std::copy(src_ptr + i * MxN, + src_ptr + (i + 1) * MxN, + dst_ptr + i * num_groups * MxN); } } } |