aboutsummaryrefslogtreecommitdiff
path: root/src/core/CL/cl_kernels/activation_layer.cl
diff options
context:
space:
mode:
authorAnthony Barbier <anthony.barbier@arm.com>2017-07-03 17:39:37 +0100
committerAnthony Barbier <anthony.barbier@arm.com>2018-09-17 14:15:39 +0100
commitac69aa137e360340fe9f148f019d93af6c3d8336 (patch)
tree88c2fb6ea8693c69b19d3b7b38fe26cd916303c8 /src/core/CL/cl_kernels/activation_layer.cl
parent05da6dd102c988081c7d5eccb227f559f740ceef (diff)
downloadComputeLibrary-ac69aa137e360340fe9f148f019d93af6c3d8336.tar.gz
COMPMID-418 Add check and fix comments after preprocessor conditions
Change-Id: I1353fd652ee180e3931e58b4ce13d651a48c7e2c Reviewed-on: http://mpd-gerrit.cambridge.arm.com/79567 Tested-by: Kaizen <jeremy.johnson+kaizengerrit@arm.com> Reviewed-by: Moritz Pflanzer <moritz.pflanzer@arm.com>
Diffstat (limited to 'src/core/CL/cl_kernels/activation_layer.cl')
-rw-r--r--src/core/CL/cl_kernels/activation_layer.cl36
1 files changed, 18 insertions, 18 deletions
diff --git a/src/core/CL/cl_kernels/activation_layer.cl b/src/core/CL/cl_kernels/activation_layer.cl
index 136191aa22..721c43c017 100644
--- a/src/core/CL/cl_kernels/activation_layer.cl
+++ b/src/core/CL/cl_kernels/activation_layer.cl
@@ -51,48 +51,48 @@
*/
__kernel void activation_layer(
TENSOR3D_DECLARATION(input)
-#if !defined IN_PLACE
+#ifndef IN_PLACE
,
TENSOR3D_DECLARATION(output)
-#endif
+#endif /* not IN_PLACE */
)
{
// Get pixels pointer
Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT(input);
-#if defined IN_PLACE
+#ifdef IN_PLACE
Tensor3D output = input;
-#else
+#else /* IN_PLACE */
Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT(output);
-#endif
+#endif /* IN_PLACE */
// Load data
VEC_DATA_TYPE(DATA_TYPE, 16)
data = vload16(0, (__global DATA_TYPE *)input.ptr);
// Perform activation
-#if defined LOGISTIC
+#ifdef LOGISTIC
data = 1 / (1 + exp(-data));
-#elif defined TANH
+#elif defined(TANH)
data = (VEC_DATA_TYPE(DATA_TYPE, 16))A * tanh((VEC_DATA_TYPE(DATA_TYPE, 16))B * data);
-#elif defined RELU
+#elif defined(RELU)
data = max(0, data);
-#elif defined BRELU
+#elif defined(BRELU)
data = min((VEC_DATA_TYPE(DATA_TYPE, 16))A, max(0, data));
-#elif defined SRELU
+#elif defined(SRELU)
data = log(1 + exp(data));
-#elif defined ABS
-#if defined TYPE_INT
+#elif defined(ABS)
+#ifdef TYPE_INT
data = abs(data);
-#else
+#else /* TYPE_INT */
data = fabs(data);
-#endif
-#elif defined SQUARE
+#endif /* TYPE_INT */
+#elif defined(SQUARE)
data = data * data;
-#elif defined SQRT
+#elif defined(SQRT)
data = sqrt(data);
-#elif defined LINEAR
+#elif defined(LINEAR)
data = (VEC_DATA_TYPE(DATA_TYPE, 16))A * data + (VEC_DATA_TYPE(DATA_TYPE, 16))B;
-#endif
+#endif /* switch TANH, RELU, BRELU, SRELU, ABS, SQUARE, SQRT, LINEAR */
// Store result
vstore16(data, 0, (__global DATA_TYPE *)output.ptr);