blob: 5f812cf5b347fadc766d250d4de1a830850510df [file] [log] [blame]
Anthony Barbier6ff3b192017-09-04 18:44:23 +01001/*
2 * Copyright (c) 2016, 2017 ARM Limited.
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
Georgios Pinitas00394ae2017-06-22 18:13:55 +010026#define TYPE VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)
27
28#if defined(FIXED_POINT_POSITION)
29#include "fixed_point.h"
30
31#define CONST_ONE (1 << FIXED_POINT_POSITION)
32#define ABS_OP(a) ABS_SAT_OP_EXPAND((a), DATA_TYPE, VEC_SIZE)
33#define ADD_OP(a, b) ADD_SAT_OP_EXPAND((a), (b), DATA_TYPE, VEC_SIZE)
34#define SUB_OP(a, b) SUB_SAT_OP_EXPAND((a), (b), DATA_TYPE, VEC_SIZE)
35#define MUL_OP(a, b) MUL_SAT_OP_EXPAND((a), (b), DATA_TYPE, VEC_SIZE, FIXED_POINT_POSITION)
36#define MLA_OP(a, b, c) MLA_SAT_OP_EXPAND((a), (b), (c), DATA_TYPE, VEC_SIZE, FIXED_POINT_POSITION)
37#define DIV_OP(a, b) DIV_SAT_OP_EXPAND((a), (b), DATA_TYPE, VEC_SIZE, FIXED_POINT_POSITION)
38#define EXP_OP(a) EXP_OP_EXPAND((a), DATA_TYPE, VEC_SIZE, FIXED_POINT_POSITION)
39#define LOG_OP(a) LOG_OP_EXPAND((a), DATA_TYPE, VEC_SIZE, FIXED_POINT_POSITION)
40#define SQRT_OP(a) DIV_OP(CONST_ONE, INVSQRT_OP_EXPAND((a), DATA_TYPE, VEC_SIZE, FIXED_POINT_POSITION))
41#define TANH_OP(a) TANH_OP_EXPAND((a), DATA_TYPE, VEC_SIZE, FIXED_POINT_POSITION)
42
43#else /* FIXED_POINT_POSITION */
44
45#define CONST_ONE (1.f)
46#define ABS_OP(a) fabs((a))
47#define ADD_OP(a, b) ((a) + (b))
48#define SUB_OP(a, b) ((a) - (b))
49#define MUL_OP(a, b) ((a) * (b))
50#define MLA_OP(a, b, c) ((b) * (c) + (a))
51#define DIV_OP(a, b) ((a) / (b))
52#define EXP_OP(a) exp((a))
53#define LOG_OP(a) log((a))
54#define SQRT_OP(a) sqrt((a))
55#define TANH_OP(a) tanh((a))
56
57#endif /* FIXED_POINT_POSITION */
58
59// Logistic Activation
60inline TYPE logistic_op(TYPE x)
61{
62 return DIV_OP(CONST_ONE, ADD_OP(CONST_ONE, EXP_OP(-x)));
63}
64// Hyperbolic Tangent Activation
65inline TYPE tanh_op(TYPE x)
66{
67 return MUL_OP((TYPE)A_VAL, TANH_OP(MUL_OP((TYPE)B_VAL, x)));
68}
69// RELU Tangent Activation
70inline TYPE relu_op(TYPE x)
71{
72 return max(0, x);
73}
74// Bounded RELU Activation
75inline TYPE brelu_op(TYPE x)
76{
77 return min((TYPE)A_VAL, max(0, x));
78}
79// Soft RELU Activation
80inline TYPE srelu_op(TYPE x)
81{
82 return LOG_OP(ADD_OP(CONST_ONE, EXP_OP(x)));
83}
84// Absolute Activation
85inline TYPE abs_op(TYPE x)
86{
87 return ABS_OP(x);
88}
89// Square Activation
90inline TYPE square_op(TYPE x)
91{
92 return MUL_OP(x, x);
93}
94// Square-root Activation
95inline TYPE sqrt_op(TYPE x)
96{
97 return SQRT_OP(x);
98}
99// Linear Activation
100inline TYPE linear_op(TYPE x)
101{
102 return MLA_OP((TYPE)B_VAL, (TYPE)A_VAL, x);
103}
104
105#define ACTIVATION_OP2(op, x) op##_op(x)
106#define ACTIVATION_OP(op, x) ACTIVATION_OP2(op, x)
107
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100108/** This performs an activation function floating point inputs.
109 *
Gian Marco Iodiceb30dcc52017-06-20 09:07:21 +0100110 * @note In order to perform the activation function "in-place", the pre-processor -DIN_PLACE must be passed at compile time
111 *
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100112 * @note Datatype should be given as a preprocessor argument using -DDATA_TYPE=type. e.g. -DDATA_TYPE=short
Georgios Pinitas00394ae2017-06-22 18:13:55 +0100113 * @note Vector size should be given as a preprocessor argument using -DVEC_SIZE=size. e.g. -DVEC_SIZE=16
114 * @note Activation function should be given as a preprocessor argument using -DACT=name. e.g. -DACT=TANH
115 * @note A, B variables required by some activation functions are set using -DA_VAL= and -DB_VAL= respectively.
116 * @note In case of fixed point calculations the fixed point position is passed using -DFIXED_POINT_POSITION=position. e.g. -DFIXED_POINT_POSITION=3.
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100117 *
Georgios Pinitas00394ae2017-06-22 18:13:55 +0100118 * @param[in] input_ptr Pointer to the source image. Supported data types: QS8/QS16/F16/F32
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100119 * @param[in] input_stride_x Stride of the source image in X dimension (in bytes)
120 * @param[in] input_step_x input_stride_x * number of elements along X processed per workitem(in bytes)
121 * @param[in] input_stride_y Stride of the source image in Y dimension (in bytes)
122 * @param[in] input_step_y input_stride_y * number of elements along Y processed per workitem(in bytes)
123 * @param[in] input_stride_z Stride of the source tensor in Z dimension (in bytes)
124 * @param[in] input_step_z input_stride_z * number of elements along Z processed per workitem(in bytes)
125 * @param[in] input_offset_first_element_in_bytes The offset of the first element in the source image
Georgios Pinitas00394ae2017-06-22 18:13:55 +0100126 * @param[out] output_ptr Pointer to the destination image. Supported data types: same as @p input_ptr
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100127 * @param[in] output_stride_x Stride of the destination image in X dimension (in bytes)
128 * @param[in] output_step_x output_stride_x * number of elements along X processed per workitem(in bytes)
129 * @param[in] output_stride_y Stride of the destination image in Y dimension (in bytes)
130 * @param[in] output_step_y output_stride_y * number of elements along Y processed per workitem(in bytes)
131 * @param[in] output_stride_z Stride of the source tensor in Z dimension (in bytes)
132 * @param[in] output_step_z output_stride_z * number of elements along Z processed per workitem(in bytes)
133 * @param[in] output_offset_first_element_in_bytes The offset of the first element in the destination image
134 */
135__kernel void activation_layer(
Gian Marco Iodiceb30dcc52017-06-20 09:07:21 +0100136 TENSOR3D_DECLARATION(input)
Anthony Barbierac69aa12017-07-03 17:39:37 +0100137#ifndef IN_PLACE
Gian Marco Iodiceb30dcc52017-06-20 09:07:21 +0100138 ,
139 TENSOR3D_DECLARATION(output)
Anthony Barbierac69aa12017-07-03 17:39:37 +0100140#endif /* not IN_PLACE */
Gian Marco Iodiceb30dcc52017-06-20 09:07:21 +0100141)
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100142{
143 // Get pixels pointer
Gian Marco Iodiceb30dcc52017-06-20 09:07:21 +0100144 Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT(input);
Anthony Barbierac69aa12017-07-03 17:39:37 +0100145#ifdef IN_PLACE
Gian Marco Iodiceb30dcc52017-06-20 09:07:21 +0100146 Tensor3D output = input;
Anthony Barbierac69aa12017-07-03 17:39:37 +0100147#else /* IN_PLACE */
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100148 Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT(output);
Anthony Barbierac69aa12017-07-03 17:39:37 +0100149#endif /* IN_PLACE */
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100150
151 // Load data
Georgios Pinitas00394ae2017-06-22 18:13:55 +0100152 TYPE data = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)input.ptr);
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100153
154 // Perform activation
Georgios Pinitas00394ae2017-06-22 18:13:55 +0100155 data = ACTIVATION_OP(ACT, data);
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100156
157 // Store result
Georgios Pinitas00394ae2017-06-22 18:13:55 +0100158 VSTORE(VEC_SIZE)
159 (data, 0, (__global DATA_TYPE *)output.ptr);
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100160}