diff options
| author | Michele Di Giorgio <michele.digiorgio@arm.com> | 2018-05-16 00:02:35 +0100 |
|---|---|---|
| committer | Anthony Barbier <anthony.barbier@arm.com> | 2018-11-02 16:52:35 +0000 |
| commit | efac7c63c045ea8fec17df54231f0d44611c7c78 (patch) | |
| tree | 74a7d16ae519a4449ebf9757199b1d5682afd0c8 | |
| parent | b5cc95bac9e3253214c56009b60a22f7e7485758 (diff) | |
| download | ComputeLibrary-efac7c63c045ea8fec17df54231f0d44611c7c78.tar.gz | |
COMPMID-959: Fix CLChannelShuffle
Change-Id: I1ea4db4e1ba37a736445ba991eeb08c247a6a61e
Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/131393
Reviewed-by: Gian Marco Iodice <gianmarco.iodice@arm.com>
Tested-by: Jenkins <bsgcomp@arm.com>
| -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); } } } |