blob: b1d5e61476d89f9d17530c819030b8ecfb262960 [file] [log] [blame]
Michalis Spyrou780db4e2017-11-23 09:49:51 +00001/*
Michele Di Giorgiod9eaf612020-07-08 11:12:57 +01002 * Copyright (c) 2017-2020 Arm Limited.
Michalis Spyrou780db4e2017-11-23 09:49:51 +00003 *
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
26/** This function applies upsample on an input image.
27 *
Michele Di Giorgiocbbed282019-12-20 13:26:08 +000028 * @param[in] src_ptr Pointer to the source image. Supported data types: All.
Michalis Spyrou780db4e2017-11-23 09:49:51 +000029 * @param[in] src_stride_x Stride of the source image in X dimension (in bytes)
30 * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
31 * @param[in] src_stride_y Stride of the source image in Y dimension (in bytes)
32 * @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
Georgios Pinitasbc06f992018-05-18 17:11:16 +010033 * @param[in] src_stride_z Stride of the source tensor in Z dimension (in bytes)
34 * @param[in] src_step_z src_stride_z * number of elements along Z processed per workitem(in bytes)
Michalis Spyrou780db4e2017-11-23 09:49:51 +000035 * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source image
Michele Di Giorgio9fef38a2018-07-06 18:06:58 +010036 * @param[out] dst_ptr Pointer to the destination image. Supported data types: same as @p src_ptr
Michalis Spyrou780db4e2017-11-23 09:49:51 +000037 * @param[in] dst_stride_x Stride of the destination image in X dimension (in bytes)
38 * @param[in] dst_step_x dst_stride_x * number of elements along X processed per workitem(in bytes)
39 * @param[in] dst_stride_y Stride of the destination image in Y dimension (in bytes)
40 * @param[in] dst_step_y dst_stride_y * number of elements along Y processed per workitem(in bytes)
Georgios Pinitasbc06f992018-05-18 17:11:16 +010041 * @param[in] dst_stride_z Stride of the source tensor in Z dimension (in bytes)
42 * @param[in] dst_step_z dst_stride_z * number of elements along Z processed per workitem(in bytes)
Michalis Spyrou780db4e2017-11-23 09:49:51 +000043 * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination image
44 */
45__kernel void deconvolution_upsample(
Georgios Pinitasbc06f992018-05-18 17:11:16 +010046 TENSOR3D_DECLARATION(src),
47 TENSOR3D_DECLARATION(dst))
Michalis Spyrou780db4e2017-11-23 09:49:51 +000048{
Georgios Pinitasbc06f992018-05-18 17:11:16 +010049 Tensor3D src = CONVERT_TO_TENSOR3D_STRUCT(src);
50 Tensor3D dst = CONVERT_TO_TENSOR3D_STRUCT(dst);
Michalis Spyrou780db4e2017-11-23 09:49:51 +000051
52 // Store result
Georgios Pinitas793f87d2018-05-18 20:08:58 +010053 *((__global DATA_TYPE *)dst.ptr) = *((__global DATA_TYPE *)src.ptr);
Michalis Spyrou780db4e2017-11-23 09:49:51 +000054}
giuros0146a49a02019-04-01 13:50:22 +010055
56#if defined(FILTER_WIDTH) && defined(FILTER_HEIGHT) && defined(SRC_WIDTH) && defined(SRC_HEIGHT) && defined(DATA_TYPE)
57/** This kernel reshapes the deconvolution output tensor before returning the result of the Deconvolution. The decovnolution output tensor
58 * is the result of a @ref CLGEMM operation between the deconvolution input and the deconvolution filter
59 *
60 * @note Data type should be given as a preprocessor argument using -DDATA_TYPE=type, e.g., -DDATA_TYPE=F32
61 * @note The width of the filter should be given as a preprocessor argument using -DFILTER_WIDTH=width, e.g., -DFILTER_WIDTH=2
62 * @note The height of the filter should be given as a preprocessor argument using -DFILTER_HEIGHT=height, e.g., -DFILTER_HEIGHT=2
63 * @note The width of the input should be given as a preprocessor argument using -DSRC_WIDTH=width, e.g., -DSRC_WIDTH=10
64 * @note The height of the input should be given as a preprocessor argument using -DSRC_HEIGHT=width, e.g., -DSRC_HEIGHT=10
65 * @note The output data layout is NHWC if the preprocessor argument NUM_FILTERS is defined, NCHW if NUM_FILTERS is not defined
66 *
Sheri Zhanga14817a2020-02-26 10:30:15 +000067 * @param[in] src_ptr Pointer to the source image. Supported data types: QASYMM8/QASYMM8_SIGNED/F16/F32
giuros0146a49a02019-04-01 13:50:22 +010068 * @param[in] src_stride_x Stride of the source image in X dimension (in bytes)
69 * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
70 * @param[in] src_stride_y Stride of the source image in Y dimension (in bytes)
71 * @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
72 * @param[in] src_stride_z Stride of the source tensor in Z dimension (in bytes)
73 * @param[in] src_step_z src_stride_z * number of elements along Z processed per workitem(in bytes)
74 * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source image
75 * @param[out] dst_ptr Pointer to the destination image. Supported data types: same as @p src_ptr
76 * @param[in] dst_stride_x Stride of the destination image in X dimension (in bytes)
77 * @param[in] dst_step_x dst_stride_x * number of elements along X processed per workitem(in bytes)
78 * @param[in] dst_stride_y Stride of the destination image in Y dimension (in bytes)
79 * @param[in] dst_step_y dst_stride_y * number of elements along Y processed per workitem(in bytes)
80 * @param[in] dst_stride_z Stride of the source tensor in Z dimension (in bytes)
81 * @param[in] dst_step_z dst_stride_z * number of elements along Z processed per workitem(in bytes)
82 * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination image
83 * @param[in] bias_ptr (Optional) Pointer to the biases vector. Supported data types: F16/F32/S32
84 * @param[in] bias_stride_x (Optional) Stride of the biases vector in X dimension (in bytes)
85 * @param[in] bias_step_x (Optional) biases_stride_x * number of elements along X processed per workitem(in bytes)
86 * @param[in] bias_offset_first_element_in_bytes (Optional) The offset of the first element in the biases vector
87 */
88__kernel void deconvolution_reshape(
89 TENSOR3D_DECLARATION(src),
90 TENSOR3D_DECLARATION(dst)
91#if defined(ADD_BIAS)
92 ,
93 VECTOR_DECLARATION(bias)
94#endif // defined(ADD_BIAS)
95)
96{
97#define FILTER_AREA ((FILTER_WIDTH) * (FILTER_HEIGHT))
98
99 Tensor3D src = CONVERT_TO_TENSOR3D_STRUCT(src);
100 Tensor3D dst = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(dst);
101 const DATA_TYPE data = *(__global DATA_TYPE *)src.ptr;
102
103 // Store result
104 const int x_in = get_global_id(0);
105 const int y_in = get_global_id(1);
106 const int z_in = get_global_id(2);
107
108#if defined(NUM_FILTERS)
109 const int bias_index = x_in / (FILTER_AREA);
110 const int z_out = bias_index + (NUM_FILTERS) * (z_in / (SRC_HEIGHT));
111 const int x_out = x_in % (FILTER_WIDTH) + y_in * (FILTER_WIDTH);
112 const int y_out = (FILTER_HEIGHT) * (z_in % (SRC_HEIGHT)) + ((x_in % (FILTER_AREA)) / (FILTER_WIDTH));
113#else // defined(NUM_FILTERS)
114 const int x_out = x_in / (FILTER_AREA);
115 const int y_out = x_in % (FILTER_WIDTH) + y_in * (FILTER_WIDTH);
116 const int z_out = (FILTER_HEIGHT) * z_in + ((x_in % (FILTER_AREA)) / (FILTER_WIDTH));
117 const int bias_index = x_out;
118#endif // defined(NUM_FILTERS)
119
120#if defined(ADD_BIAS)
121 Vector bias = CONVERT_TO_VECTOR_STRUCT_NO_STEP(bias);
122 const DATA_TYPE bias_val = *(__global DATA_TYPE *)vector_offset(&bias, bias_index);
123 *((__global DATA_TYPE *)tensor3D_offset(&dst, x_out, y_out, z_out)) = data + bias_val;
124#else // defined(ADD_BIAS)
125 *((__global DATA_TYPE *)tensor3D_offset(&dst, x_out, y_out, z_out)) = data;
126#endif // defined(ADD_BIAS)
127
128#undef FILTER_AREA
129}
130#endif // defined(FILTER_WIDTH) && defined(FILTER_HEIGHT) && defined(SRC_WIDTH) && defined(SRC_HEIGHT) && defined(DATA_TYPE)