blob: d44e78d990ff0c4c63bdc0949c06907a6dd9d645 [file] [log] [blame]
Georgios Pinitas5ee66ea2017-09-07 17:29:16 +01001/*
Adnan AlSinan7075fe22021-07-05 13:12:52 +01002 * Copyright (c) 2019-2021 Arm Limited.
Georgios Pinitas5ee66ea2017-09-07 17:29:16 +01003 *
4 * SPDX-License-Identifier: MIT
5 *
6 * Permission is hereby granted, free of charge, to any person obtaining a copy
7 * of this software and associated documentation files (the "Software"), to
8 * deal in the Software without restriction, including without limitation the
9 * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
10 * sell copies of the Software, and to permit persons to whom the Software is
11 * furnished to do so, subject to the following conditions:
12 *
13 * The above copyright notice and this permission notice shall be included in all
14 * copies or substantial portions of the Software.
15 *
16 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
17 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
18 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
19 * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
20 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
21 * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
22 * SOFTWARE.
23 */
24#include "helpers.h"
25
Adnan AlSinan7075fe22021-07-05 13:12:52 +010026#if defined(DATA_TYPE) && defined(BLOCK_SHAPE) && defined(CHANNEL_SIZE)
27/** Space to depth transformation. (NHWC)
Georgios Pinitas5ee66ea2017-09-07 17:29:16 +010028 *
Adnan AlSinan7075fe22021-07-05 13:12:52 +010029 * @note Datatype should be given as a preprocessor argument using -DDATA_TYPE=type. e.g. -DDATA_TYPE=float
30 * @note The input tensor batch size must be passed at compile time using -DCHANNEL_SIZE. e.g. -DCHANNEL_SIZE=2
31 * @note The block shape must be passed at compile time using -DBLOCK_SHAPE. e.g. -DBLOCK_SHAPE=2
Georgios Pinitas5ee66ea2017-09-07 17:29:16 +010032 *
Adnan AlSinan7075fe22021-07-05 13:12:52 +010033 * @param[in] input_ptr Pointer to the source tensor. Supported data types: All
34 * @param[in] input_stride_x Stride of the source tensor in X dimension (in bytes)
Georgios Pinitas5ee66ea2017-09-07 17:29:16 +010035 * @param[in] input_step_x input_stride_x * number of elements along X processed per workitem(in bytes)
Adnan AlSinan7075fe22021-07-05 13:12:52 +010036 * @param[in] input_stride_y Stride of the source tensor in Y dimension (in bytes)
Georgios Pinitas5ee66ea2017-09-07 17:29:16 +010037 * @param[in] input_step_y input_stride_y * number of elements along Y processed per workitem(in bytes)
Adnan AlSinan7075fe22021-07-05 13:12:52 +010038 * @param[in] input_stride_z Stride of the source tensor in Z dimension (in bytes)
Georgios Pinitas5ee66ea2017-09-07 17:29:16 +010039 * @param[in] input_step_z input_stride_z * number of elements along Z processed per workitem(in bytes)
40 * @param[in] input_offset_first_element_in_bytes The offset of the first element in the first source tensor
Adnan AlSinan7075fe22021-07-05 13:12:52 +010041 * @param[in] batch_id The input tensor batch id
Georgios Pinitas5ee66ea2017-09-07 17:29:16 +010042 * @param[out] output_ptr Pointer to the destination tensor. Supported data types: same as @p input_ptr
43 * @param[in] output_stride_x Stride of the destination tensor in X dimension (in bytes)
44 * @param[in] output_step_x output_stride_x * number of elements along X processed per workitem(in bytes)
45 * @param[in] output_stride_y Stride of the destination tensor in Y dimension (in bytes)
46 * @param[in] output_step_y output_stride_y * number of elements along Y processed per workitem(in bytes)
Adnan AlSinan7075fe22021-07-05 13:12:52 +010047 * @param[in] output_stride_z Stride of the source tensor in Z dimension (in bytes)
Georgios Pinitas5ee66ea2017-09-07 17:29:16 +010048 * @param[in] output_step_z output_stride_z * number of elements along Z processed per workitem(in bytes)
49 * @param[in] output_offset_first_element_in_bytes The offset of the first element in the destination tensor
Georgios Pinitas5ee66ea2017-09-07 17:29:16 +010050 */
Adnan AlSinan7075fe22021-07-05 13:12:52 +010051__kernel void space_to_depth_nhwc(
52 TENSOR4D_DECLARATION(input),
53 const int batch_id,
54 TENSOR3D_DECLARATION(output))
Georgios Pinitas5ee66ea2017-09-07 17:29:16 +010055{
Adnan AlSinan7075fe22021-07-05 13:12:52 +010056 Tensor4D in = CONVERT_TO_TENSOR4D_STRUCT_NO_STEP(input, 0);
57 Tensor3D out = CONVERT_TO_TENSOR3D_STRUCT(output);
Georgios Pinitas5ee66ea2017-09-07 17:29:16 +010058
Adnan AlSinan7075fe22021-07-05 13:12:52 +010059 const int r = (CHANNEL_SIZE / (BLOCK_SHAPE * BLOCK_SHAPE));
60 const int x = get_global_id(1);
61 const int y = get_global_id(2);
62 const int z = get_global_id(0) % r;
Georgios Pinitas5ee66ea2017-09-07 17:29:16 +010063
Adnan AlSinan7075fe22021-07-05 13:12:52 +010064 const int in_x = x * BLOCK_SHAPE + (get_global_id(0) / r) % BLOCK_SHAPE;
65 const int in_y = y * BLOCK_SHAPE + (get_global_id(0) / r) / BLOCK_SHAPE;
Georgios Pinitas5ee66ea2017-09-07 17:29:16 +010066
Adnan AlSinan7075fe22021-07-05 13:12:52 +010067 *((__global DATA_TYPE *)out.ptr) = *((__global DATA_TYPE *)tensor4D_offset(&in, z, in_x, in_y, batch_id));
Georgios Pinitas5ee66ea2017-09-07 17:29:16 +010068}
Adnan AlSinan7075fe22021-07-05 13:12:52 +010069#endif // defined(DATA_TYPE) && defined(BLOCK_SHAPE) && defined(CHANNEL_SIZE)