Gitiles
Code Review Sign In
review.mlplatform.org / ml / ComputeLibrary / be4100605230868b5cc50cabee395613dbfb62cd / . / src / core / CL / cl_kernels / convolution_layer.cl
blob: 2b83e5adf18cc3accb5a3b7f2e6ee37af3872772 [file] [log] [blame]
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]1/*
Gian Marco76faef82018-01-29 12:15:32 +0000[diff] [blame]2 * Copyright (c) 2017-2018 ARM Limited.
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]3 *
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
Gian Marco76faef82018-01-29 12:15:32 +0000[diff] [blame]26#if defined(DATA_TYPE)
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]27/** This kernel reshapes the tensor's low three dimensions to single column
28 *
29 * @note Datatype should be given as a preprocessor argument using -DDATA_TYPE=type. e.g. -DDATA_TYPE=short
30 *
Gian Marco Iodice13edbff2017-06-26 17:20:16 +0100[diff] [blame]31 * @param[in] src_ptr Pointer to the source tensor. Supported data types: F16/F32
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]32 * @param[in] src_stride_x Stride of the source tensor in X dimension (in bytes)
33 * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
34 * @param[in] src_stride_y Stride of the source tensor in Y dimension (in bytes)
35 * @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
36 * @param[in] src_stride_z Stride of the source tensor in Z dimension (in bytes)
37 * @param[in] src_step_z src_stride_z * number of elements along Y processed per workitem(in bytes)
38 * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source tensor
Gian Marco Iodice13edbff2017-06-26 17:20:16 +0100[diff] [blame]39 * @param[out] dst_ptr Pointer to the destination tensor. Same as @p src_ptr
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]40 * @param[in] dst_stride_x Stride of the destination tensor in X dimension (in bytes)
41 * @param[in] dst_step_x dst_stride_x * number of elements along X processed per workitem(in bytes)
42 * @param[in] dst_stride_y Stride of the destination tensor in Y dimension (in bytes)
43 * @param[in] dst_step_y dst_stride_y * number of elements along Y processed per workitem(in bytes)
44 * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination tensor
Gian Marco Iodice13edbff2017-06-26 17:20:16 +0100[diff] [blame]45 * @param[in] bias_ptr Pointer to the bias tensor. Same as @p src_ptr
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]46 * @param[in] bias_stride_x Stride of the bias tensor in X dimension (in bytes)
47 * @param[in] bias_step_x bias_stride_x * number of elements along X processed per workitem(in bytes)
48 * @param[in] bias_offset_first_element_in_bytes The offset of the first element in the source tensor
49 * @param[in] width The width of the input tensor
50 * @param[in] height The height of the input tensor
51 * @param[in] depth The depth of the input tensor
52 * @param[in] total_filters Total number of filters. 4th dimension of the weights matrix
53 */
Georgios Pinitas19ea4192018-06-19 13:09:53 +0100[diff] [blame]54__kernel void reshape_to_columns_nchw(
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]55 TENSOR3D_DECLARATION(src),
56 IMAGE_DECLARATION(dst),
Anthony Barbierac69aa12017-07-03 17:39:37 +0100[diff] [blame]57#ifdef HAS_BIAS
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]58 VECTOR_DECLARATION(bias),
Anthony Barbierac69aa12017-07-03 17:39:37 +0100[diff] [blame]59#endif /* HAS_BIAS */
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]60 uint width, uint height, uint depth, uint total_filters)
61{
62 Tensor3D src = CONVERT_TO_TENSOR3D_STRUCT(src);
63 bool is_last_thread = (get_global_id(0) == (get_global_size(0) - 1) && get_global_id(1) == (get_global_size(1) - 1) && get_global_id(2) == (get_global_size(2) - 1));
64
65 __global uchar *tmp_src_ptr = src.ptr;
66 __global uchar *tmp_dst_ptr = dst_ptr + dst_offset_first_element_in_bytes + get_global_id(0) * dst_stride_y + get_global_id(1) * width * dst_stride_y + get_global_id(
67 2) * width * height * dst_stride_y;
Anthony Barbierac69aa12017-07-03 17:39:37 +0100[diff] [blame]68#ifdef HAS_BIAS
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]69 __global uchar *tmp_bias_ptr = bias_ptr + bias_offset_first_element_in_bytes;
Anthony Barbierac69aa12017-07-03 17:39:37 +0100[diff] [blame]70#endif /* HAS_BIAS */
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]71
72 if(is_last_thread)
73 {
74 for(uint i = 0; i < total_filters; ++i)
75 {
76 *((__global DATA_TYPE *)tmp_dst_ptr) = *((__global DATA_TYPE *)tmp_src_ptr);
77
Anthony Barbierac69aa12017-07-03 17:39:37 +0100[diff] [blame]78#ifdef HAS_BIAS
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]79 *((__global DATA_TYPE *)(tmp_dst_ptr + dst_stride_y)) = *((__global DATA_TYPE *)(tmp_bias_ptr));
80 tmp_bias_ptr += bias_stride_x;
Anthony Barbierac69aa12017-07-03 17:39:37 +0100[diff] [blame]81#endif /* HAS_BIAS */
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]82 tmp_src_ptr += depth * src_stride_z;
83 tmp_dst_ptr += dst_stride_x;
84 }
85 }
86 else
87 {
88 for(uint i = 0; i < total_filters; ++i)
89 {
90 *((__global DATA_TYPE *)tmp_dst_ptr) = *((__global DATA_TYPE *)tmp_src_ptr);
91 tmp_src_ptr += depth * src_stride_z;
92 tmp_dst_ptr += dst_stride_x;
93 }
94 }
95}
Georgios Pinitas19ea4192018-06-19 13:09:53 +0100[diff] [blame]96
97/** This kernel reshapes the tensor's low three dimensions to single column
98 *
99 * @note Datatype should be given as a preprocessor argument using -DDATA_TYPE=type. e.g. -DDATA_TYPE=short
100 *
101 * @param[in] src_ptr Pointer to the source tensor. Supported data types: F16/F32
102 * @param[in] src_stride_x Stride of the source tensor in X dimension (in bytes)
103 * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
104 * @param[in] src_stride_y Stride of the source tensor in Y dimension (in bytes)
105 * @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
106 * @param[in] src_stride_z Stride of the source tensor in Z dimension (in bytes)
107 * @param[in] src_step_z src_stride_z * number of elements along Y processed per workitem(in bytes)
108 * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source tensor
109 * @param[out] dst_ptr Pointer to the destination tensor. Same as @p src_ptr
110 * @param[in] dst_stride_x Stride of the destination tensor in X dimension (in bytes)
111 * @param[in] dst_step_x dst_stride_x * number of elements along X processed per workitem(in bytes)
112 * @param[in] dst_stride_y Stride of the destination tensor in Y dimension (in bytes)
113 * @param[in] dst_step_y dst_stride_y * number of elements along Y processed per workitem(in bytes)
114 * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination tensor
115 * @param[in] bias_ptr Pointer to the bias tensor. Same as @p src_ptr
116 * @param[in] bias_stride_x Stride of the bias tensor in X dimension (in bytes)
117 * @param[in] bias_step_x bias_stride_x * number of elements along X processed per workitem(in bytes)
118 * @param[in] bias_offset_first_element_in_bytes The offset of the first element in the source tensor
119 * @param[in] depth The depth of the input tensor
120 * @param[in] width The width of the input tensor
121 * @param[in] height The height of the input tensor
122 * @param[in] total_filters Total number of filters. 4th dimension of the weights matrix
123 */
124__kernel void reshape_to_columns_nhwc(
125 TENSOR3D_DECLARATION(src),
126 IMAGE_DECLARATION(dst),
127#ifdef HAS_BIAS
128 VECTOR_DECLARATION(bias),
129#endif /* HAS_BIAS */
130 uint depth, uint width, uint height, uint total_filters)
131{
132 Tensor3D src = CONVERT_TO_TENSOR3D_STRUCT(src);
133 bool is_last_thread = (get_global_id(0) == (get_global_size(0) - 1) && get_global_id(1) == (get_global_size(1) - 1) && get_global_id(2) == (get_global_size(2) - 1));
134
135 __global uchar *tmp_src_ptr = src.ptr;
136 __global uchar *tmp_dst_ptr = dst_ptr + dst_offset_first_element_in_bytes + get_global_id(1) * dst_stride_y + get_global_id(2) * width * dst_stride_y + get_global_id(
137 0) * width * height * dst_stride_y;
138#ifdef HAS_BIAS
139 __global uchar *tmp_bias_ptr = bias_ptr + bias_offset_first_element_in_bytes;
140#endif /* HAS_BIAS */
141
142 if(is_last_thread)
143 {
144 for(uint i = 0; i < total_filters; ++i)
145 {
146 *((__global DATA_TYPE *)tmp_dst_ptr) = *((__global DATA_TYPE *)tmp_src_ptr);
147
148#ifdef HAS_BIAS
149 *((__global DATA_TYPE *)(tmp_dst_ptr + dst_stride_y)) = *((__global DATA_TYPE *)(tmp_bias_ptr));
150 tmp_bias_ptr += bias_stride_x;
151#endif /* HAS_BIAS */
152 tmp_src_ptr += height * src_stride_z;
153 tmp_dst_ptr += dst_stride_x;
154 }
155 }
156 else
157 {
158 for(uint i = 0; i < total_filters; ++i)
159 {
160 *((__global DATA_TYPE *)tmp_dst_ptr) = *((__global DATA_TYPE *)tmp_src_ptr);
161 tmp_src_ptr += height * src_stride_z;
162 tmp_dst_ptr += dst_stride_x;
163 }
164 }
165}
Gian Marco76faef82018-01-29 12:15:32 +0000[diff] [blame]166#endif // defined(DATA_TYPE)