aboutsummaryrefslogtreecommitdiff
path: root/src/core/CL/cl_kernels/activation_layer_qa8.cl
diff options
context:
space:
mode:
Diffstat (limited to 'src/core/CL/cl_kernels/activation_layer_qa8.cl')
-rw-r--r--src/core/CL/cl_kernels/activation_layer_qa8.cl13
1 files changed, 12 insertions, 1 deletions
diff --git a/src/core/CL/cl_kernels/activation_layer_qa8.cl b/src/core/CL/cl_kernels/activation_layer_qa8.cl
index 66e54ed6ad..8f6a807613 100644
--- a/src/core/CL/cl_kernels/activation_layer_qa8.cl
+++ b/src/core/CL/cl_kernels/activation_layer_qa8.cl
@@ -24,7 +24,18 @@
#include "helpers.h"
#define TYPE VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)
+#define VEC_FLOAT VEC_DATA_TYPE(float, VEC_SIZE)
+// Logistic Activation
+inline TYPE logistic_op(TYPE x)
+{
+ VEC_FLOAT x_flt = CONVERT(x, VEC_FLOAT);
+ x_flt = round(x_flt - (float)O1_VAL) * ((float)S1_VAL);
+ x_flt = 1.f / (1.f + exp(-x_flt));
+
+ const TYPE x_u8 = CONVERT_SAT(round(x_flt / ((float)S1_VAL)) + (float)O1_VAL, TYPE);
+ return x_u8;
+}
// RELU Activation
inline TYPE relu_op(TYPE x)
{
@@ -119,4 +130,4 @@ __kernel void activation_layer_qa8(
(data, 0, (__global DATA_TYPE *)output.ptr);
}
-#endif /* defined(ACT) */ \ No newline at end of file
+#endif /* defined(ACT) */