Gitiles
Code Review Sign In
review.mlplatform.org / ml / ComputeLibrary / 83e3e7537bda871c2aade349d7f584d065a75092 / . / src / core / CL / cl_kernels / arithmetic_op.cl
blob: 557615e7f2c978493399099ddf591e47977785a9 [file] [log] [blame]
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]1/*
Michalis Spyrou0a887922018-06-11 16:30:23 +0100[diff] [blame]2 * Copyright (c) 2016-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
26#ifdef SATURATE
27#define ADD(x, y) add_sat((x), (y))
28#define SUB(x, y) sub_sat((x), (y))
Anthony Barbierac69aa12017-07-03 17:39:37 +0100[diff] [blame]29#else /* SATURATE */
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]30#define ADD(x, y) (x) + (y)
31#define SUB(x, y) (x) - (y)
Anthony Barbierac69aa12017-07-03 17:39:37 +0100[diff] [blame]32#endif /* SATURATE */
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]33
Michalis Spyrou0a887922018-06-11 16:30:23 +0100[diff] [blame]34#define DIV(x, y) (x) / (y)
35
Michele Di Giorgioa1422fb2018-10-24 12:20:19 +0100[diff] [blame]36#if defined(DATA_TYPE_IN1) && defined(DATA_TYPE_IN2) && defined(DATA_TYPE_OUT) && defined(VEC_SIZE)
Georgios Pinitas1d08a312018-01-03 12:29:22 +0000[diff] [blame]37/** This function adds two tensors.
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]38 *
39 * @attention The input and output data_types need to be passed at compile time using -DDATA_TYPE_IN1, -DDATA_TYPE_IN2 and -DDATA_TYPE_OUT:
40 * e.g. -DDATA_TYPE_IN1=uchar -DDATA_TYPE_IN2=uchar -DDATA_TYPE_OUT=short
41 * @attention To perform saturating operation -DSATURATE has to be passed to the compiler otherwise wrapping policy will be used.
Michele Di Giorgioa1422fb2018-10-24 12:20:19 +0100[diff] [blame]42 * @attention Vector size should be given as a preprocessor argument using -DVEC_SIZE=size. e.g. -DVEC_SIZE=16
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]43 *
Vidhya Sudhan Loganathan7485d5a2018-07-04 09:34:00 +0100[diff] [blame]44 * @param[in] in1_ptr Pointer to the source tensor. Supported data types: U8/S16/F16/F32
Georgios Pinitas1d08a312018-01-03 12:29:22 +0000[diff] [blame]45 * @param[in] in1_stride_x Stride of the source tensor in X dimension (in bytes)
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]46 * @param[in] in1_step_x in1_stride_x * number of elements along X processed per workitem(in bytes)
Georgios Pinitas1d08a312018-01-03 12:29:22 +0000[diff] [blame]47 * @param[in] in1_stride_y Stride of the source tensor in Y dimension (in bytes)
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]48 * @param[in] in1_step_y in1_stride_y * number of elements along Y processed per workitem(in bytes)
Georgios Pinitas1d08a312018-01-03 12:29:22 +0000[diff] [blame]49 * @param[in] in1_stride_z Stride of the source tensor in Z dimension (in bytes)
50 * @param[in] in1_step_z in1_stride_z * number of elements along Z processed per workitem(in bytes)
51 * @param[in] in1_offset_first_element_in_bytes The offset of the first element in the source tensor
Vidhya Sudhan Loganathan7485d5a2018-07-04 09:34:00 +0100[diff] [blame]52 * @param[in] in2_ptr Pointer to the source tensor. Supported data types: U8/S16/F16/F32
Georgios Pinitas1d08a312018-01-03 12:29:22 +0000[diff] [blame]53 * @param[in] in2_stride_x Stride of the source tensor in X dimension (in bytes)
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]54 * @param[in] in2_step_x in2_stride_x * number of elements along X processed per workitem(in bytes)
Georgios Pinitas1d08a312018-01-03 12:29:22 +0000[diff] [blame]55 * @param[in] in2_stride_y Stride of the source tensor in Y dimension (in bytes)
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]56 * @param[in] in2_step_y in2_stride_y * number of elements along Y processed per workitem(in bytes)
Georgios Pinitas1d08a312018-01-03 12:29:22 +0000[diff] [blame]57 * @param[in] in2_stride_z Stride of the source tensor in Z dimension (in bytes)
58 * @param[in] in2_step_z in2_stride_z * number of elements along Z processed per workitem(in bytes)
59 * @param[in] in2_offset_first_element_in_bytes The offset of the first element in the source tensor
Vidhya Sudhan Loganathan7485d5a2018-07-04 09:34:00 +0100[diff] [blame]60 * @param[out] out_ptr Pointer to the destination tensor. Supported data types: U8 (only if both inputs are U8), S16/F16/F32
Georgios Pinitas1d08a312018-01-03 12:29:22 +0000[diff] [blame]61 * @param[in] out_stride_x Stride of the destination tensor in X dimension (in bytes)
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]62 * @param[in] out_step_x out_stride_x * number of elements along X processed per workitem(in bytes)
Georgios Pinitas1d08a312018-01-03 12:29:22 +0000[diff] [blame]63 * @param[in] out_stride_y Stride of the destination tensor in Y dimension (in bytes)
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]64 * @param[in] out_step_y out_stride_y * number of elements along Y processed per workitem(in bytes)
Georgios Pinitas1d08a312018-01-03 12:29:22 +0000[diff] [blame]65 * @param[in] out_stride_z Stride of the source tensor in Z dimension (in bytes)
66 * @param[in] out_step_z out_stride_z * number of elements along Z processed per workitem(in bytes)
67 * @param[in] out_offset_first_element_in_bytes The offset of the first element in the destination tensor
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]68 */
69__kernel void arithmetic_add(
Georgios Pinitas1d08a312018-01-03 12:29:22 +0000[diff] [blame]70 TENSOR3D_DECLARATION(in1),
71 TENSOR3D_DECLARATION(in2),
72 TENSOR3D_DECLARATION(out))
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]73{
74 // Get pixels pointer
Georgios Pinitas1d08a312018-01-03 12:29:22 +0000[diff] [blame]75 Tensor3D in1 = CONVERT_TO_TENSOR3D_STRUCT(in1);
76 Tensor3D in2 = CONVERT_TO_TENSOR3D_STRUCT(in2);
77 Tensor3D out = CONVERT_TO_TENSOR3D_STRUCT(out);
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]78
79 // Load values
Michele Di Giorgioa1422fb2018-10-24 12:20:19 +0100[diff] [blame]80 VEC_DATA_TYPE(DATA_TYPE_OUT, VEC_SIZE)
81 in_a = CONVERT(VLOAD(VEC_SIZE)(0, (__global DATA_TYPE_IN1 *)in1.ptr), VEC_DATA_TYPE(DATA_TYPE_OUT, VEC_SIZE));
82 VEC_DATA_TYPE(DATA_TYPE_OUT, VEC_SIZE)
83 in_b = CONVERT(VLOAD(VEC_SIZE)(0, (__global DATA_TYPE_IN2 *)in2.ptr), VEC_DATA_TYPE(DATA_TYPE_OUT, VEC_SIZE));
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]84
85 // Calculate and store result
Michele Di Giorgioa1422fb2018-10-24 12:20:19 +0100[diff] [blame]86 VSTORE(VEC_SIZE)
87 (ADD(in_a, in_b), 0, (__global DATA_TYPE_OUT *)out.ptr);
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]88}
Michele Di Giorgioa1422fb2018-10-24 12:20:19 +0100[diff] [blame]89#endif /* defined(DATA_TYPE_IN1) && defined(DATA_TYPE_IN2) && defined(DATA_TYPE_OUT) && defined(VEC_SIZE) */
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]90
Michalis Spyrou0a887922018-06-11 16:30:23 +0100[diff] [blame]91/** This function subtracts one tensor from another.
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]92 *
93 * @attention The input and output data_types need to be passed at compile time using -DDATA_TYPE_IN1, -DDATA_TYPE_IN2 and -DDATA_TYPE_OUT:
94 * e.g. -DDATA_TYPE_IN1=uchar -DDATA_TYPE_IN2=uchar -DDATA_TYPE_OUT=short
95 * @attention To perform saturating operation -DSATURATE has to be passed to the compiler otherwise wrapping policy will be used.
96 *
Georgios Pinitas1d08a312018-01-03 12:29:22 +0000[diff] [blame]97 * @param[in] in1_ptr Pointer to the source tensor. Supported data types: U8, S16
98 * @param[in] in1_stride_x Stride of the source tensor in X dimension (in bytes)
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]99 * @param[in] in1_step_x in1_stride_x * number of elements along X processed per workitem(in bytes)
Georgios Pinitas1d08a312018-01-03 12:29:22 +0000[diff] [blame]100 * @param[in] in1_stride_y Stride of the source tensor in Y dimension (in bytes)
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]101 * @param[in] in1_step_y in1_stride_y * number of elements along Y processed per workitem(in bytes)
Georgios Pinitas1d08a312018-01-03 12:29:22 +0000[diff] [blame]102 * @param[in] in1_stride_z Stride of the source tensor in Z dimension (in bytes)
103 * @param[in] in1_step_z in1_stride_z * number of elements along Z processed per workitem(in bytes)
104 * @param[in] in1_offset_first_element_in_bytes The offset of the first element in the source tensor
105 * @param[in] in2_ptr Pointer to the source tensor. Supported data types: U8, S16
106 * @param[in] in2_stride_x Stride of the source tensor in X dimension (in bytes)
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]107 * @param[in] in2_step_x in2_stride_x * number of elements along X processed per workitem(in bytes)
Georgios Pinitas1d08a312018-01-03 12:29:22 +0000[diff] [blame]108 * @param[in] in2_stride_y Stride of the source tensor in Y dimension (in bytes)
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]109 * @param[in] in2_step_y in2_stride_y * number of elements along Y processed per workitem(in bytes)
Georgios Pinitas1d08a312018-01-03 12:29:22 +0000[diff] [blame]110 * @param[in] in2_stride_z Stride of the source tensor in Z dimension (in bytes)
111 * @param[in] in2_step_z in2_stride_z * number of elements along Z processed per workitem(in bytes)
112 * @param[in] in2_offset_first_element_in_bytes The offset of the first element in the source tensor
113 * @param[out] out_ptr Pointer to the destination tensor. Supported data types: U8, S16
114 * @param[in] out_stride_x Stride of the destination tensor in X dimension (in bytes)
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]115 * @param[in] out_step_x out_stride_x * number of elements along X processed per workitem(in bytes)
Georgios Pinitas1d08a312018-01-03 12:29:22 +0000[diff] [blame]116 * @param[in] out_stride_y Stride of the destination tensor in Y dimension (in bytes)
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]117 * @param[in] out_step_y out_stride_y * number of elements along Y processed per workitem(in bytes)
Georgios Pinitas1d08a312018-01-03 12:29:22 +0000[diff] [blame]118 * @param[in] out_stride_z Stride of the source tensor in Z dimension (in bytes)
119 * @param[in] out_step_z out_stride_z * number of elements along Z processed per workitem(in bytes)
120 * @param[in] out_offset_first_element_in_bytes The offset of the first element in the destination tensor
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]121 */
122__kernel void arithmetic_sub(
Georgios Pinitas1d08a312018-01-03 12:29:22 +0000[diff] [blame]123 TENSOR3D_DECLARATION(in1),
124 TENSOR3D_DECLARATION(in2),
125 TENSOR3D_DECLARATION(out))
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]126{
127 // Get pixels pointer
Georgios Pinitas1d08a312018-01-03 12:29:22 +0000[diff] [blame]128 Tensor3D in1 = CONVERT_TO_TENSOR3D_STRUCT(in1);
129 Tensor3D in2 = CONVERT_TO_TENSOR3D_STRUCT(in2);
130 Tensor3D out = CONVERT_TO_TENSOR3D_STRUCT(out);
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]131
132 // Load values
133 VEC_DATA_TYPE(DATA_TYPE_OUT, 16)
134 in_a = CONVERT(vload16(0, (__global DATA_TYPE_IN1 *)in1.ptr), VEC_DATA_TYPE(DATA_TYPE_OUT, 16));
135 VEC_DATA_TYPE(DATA_TYPE_OUT, 16)
136 in_b = CONVERT(vload16(0, (__global DATA_TYPE_IN2 *)in2.ptr), VEC_DATA_TYPE(DATA_TYPE_OUT, 16));
137
138 // Calculate and store result
139 vstore16(SUB(in_a, in_b), 0, (__global DATA_TYPE_OUT *)out.ptr);
140}
Michalis Spyrou0a887922018-06-11 16:30:23 +0100[diff] [blame]141
142/** This function divides one tensor from another.
143 *
144 * @attention The input and output data_types need to be passed at compile time using -DDATA_TYPE_IN1, -DDATA_TYPE_IN2 and -DDATA_TYPE_OUT:
145 * e.g. -DDATA_TYPE_IN1=float -DDATA_TYPE_IN2=float -DDATA_TYPE_OUT=float
146 *
147 * @param[in] in1_ptr Pointer to the source tensor. Supported data types: F16/F32
148 * @param[in] in1_stride_x Stride of the source tensor in X dimension (in bytes)
149 * @param[in] in1_step_x in1_stride_x * number of elements along X processed per workitem(in bytes)
150 * @param[in] in1_stride_y Stride of the source tensor in Y dimension (in bytes)
151 * @param[in] in1_step_y in1_stride_y * number of elements along Y processed per workitem(in bytes)
152 * @param[in] in1_stride_z Stride of the source tensor in Z dimension (in bytes)
153 * @param[in] in1_step_z in1_stride_z * number of elements along Z processed per workitem(in bytes)
154 * @param[in] in1_offset_first_element_in_bytes The offset of the first element in the source tensor
155 * @param[in] in2_ptr Pointer to the source tensor. Supported data types: Same as @p in1_ptr
156 * @param[in] in2_stride_x Stride of the source tensor in X dimension (in bytes)
157 * @param[in] in2_step_x in2_stride_x * number of elements along X processed per workitem(in bytes)
158 * @param[in] in2_stride_y Stride of the source tensor in Y dimension (in bytes)
159 * @param[in] in2_step_y in2_stride_y * number of elements along Y processed per workitem(in bytes)
160 * @param[in] in2_stride_z Stride of the source tensor in Z dimension (in bytes)
161 * @param[in] in2_step_z in2_stride_z * number of elements along Z processed per workitem(in bytes)
162 * @param[in] in2_offset_first_element_in_bytes The offset of the first element in the source tensor
163 * @param[out] out_ptr Pointer to the destination tensor. Supported data types: Same as @p in1_ptr
164 * @param[in] out_stride_x Stride of the destination tensor in X dimension (in bytes)
165 * @param[in] out_step_x out_stride_x * number of elements along X processed per workitem(in bytes)
166 * @param[in] out_stride_y Stride of the destination tensor in Y dimension (in bytes)
167 * @param[in] out_step_y out_stride_y * number of elements along Y processed per workitem(in bytes)
168 * @param[in] out_stride_z Stride of the source tensor in Z dimension (in bytes)
169 * @param[in] out_step_z out_stride_z * number of elements along Z processed per workitem(in bytes)
170 * @param[in] out_offset_first_element_in_bytes The offset of the first element in the destination tensor
171 */
172__kernel void arithmetic_div(
173 TENSOR3D_DECLARATION(in1),
174 TENSOR3D_DECLARATION(in2),
175 TENSOR3D_DECLARATION(out))
176{
177 // Get pixels pointer
178 Tensor3D in1 = CONVERT_TO_TENSOR3D_STRUCT(in1);
179 Tensor3D in2 = CONVERT_TO_TENSOR3D_STRUCT(in2);
180 Tensor3D out = CONVERT_TO_TENSOR3D_STRUCT(out);
181
182 // Load values
183 VEC_DATA_TYPE(DATA_TYPE_OUT, 16)
184 in_a = CONVERT(vload16(0, (__global DATA_TYPE_IN1 *)in1.ptr), VEC_DATA_TYPE(DATA_TYPE_OUT, 16));
185 VEC_DATA_TYPE(DATA_TYPE_OUT, 16)
186 in_b = CONVERT(vload16(0, (__global DATA_TYPE_IN2 *)in2.ptr), VEC_DATA_TYPE(DATA_TYPE_OUT, 16));
187
188 // Calculate and store result
189 vstore16(DIV(in_a, in_b), 0, (__global DATA_TYPE_OUT *)out.ptr);
190}