From d304e80112e5d246b4bb884d1211eef9381bf034 Mon Sep 17 00:00:00 2001 From: Michele Di Giorgio Date: Fri, 6 Jul 2018 10:17:33 +0100 Subject: COMPMID-1349: Add support for QASYMM8 LOGISTIC activation in CLActivationLayer Change-Id: Ibabce61cf5427de80078a6468023bed05f5e7c2c Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/139006 Tested-by: Jenkins Reviewed-by: Anthony Barbier Reviewed-by: Georgios Pinitas --- src/core/CL/cl_kernels/activation_layer_qa8.cl | 13 ++++++++++++- src/core/CL/kernels/CLActivationLayerKernel.cpp | 13 +++++++------ .../CL/kernels/CLDepthwiseConvolutionLayer3x3NCHWKernel.cpp | 5 +++-- .../CL/kernels/CLDepthwiseConvolutionLayer3x3NHWCKernel.cpp | 5 +++-- 4 files changed, 25 insertions(+), 11 deletions(-) (limited to 'src/core/CL') 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) */ diff --git a/src/core/CL/kernels/CLActivationLayerKernel.cpp b/src/core/CL/kernels/CLActivationLayerKernel.cpp index 1ae1032cba..d8bd2f7ee1 100644 --- a/src/core/CL/kernels/CLActivationLayerKernel.cpp +++ b/src/core/CL/kernels/CLActivationLayerKernel.cpp @@ -49,8 +49,9 @@ Status validate_arguments(const ITensorInfo *input, const ITensorInfo *output, c ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::U8, DataType::QASYMM8, DataType::F16, DataType::F32); ARM_COMPUTE_RETURN_ERROR_ON_MSG((input->data_type() == DataType::QASYMM8) && (act_info.activation() != ActivationLayerInfo::ActivationFunction::LU_BOUNDED_RELU) && (act_info.activation() != ActivationLayerInfo::ActivationFunction::BOUNDED_RELU) - && (act_info.activation() != ActivationLayerInfo::ActivationFunction::RELU), - "For QASYMM8 only relu, lower bounded relu and lower-upper bounded relu are supported"); + && (act_info.activation() != ActivationLayerInfo::ActivationFunction::RELU) + && (act_info.activation() != ActivationLayerInfo::ActivationFunction::LOGISTIC), + "For QASYMM8 only logistic, relu, lower bounded relu and lower-upper bounded relu are supported"); // Checks performed when output is configured if((output != nullptr) && (output->total_size() != 0)) @@ -139,22 +140,22 @@ void CLActivationLayerKernel::configure(ICLTensor *input, ICLTensor *output, Act build_opts.emplace(("-DA_VAL=" + support::cpp11::to_string(a_const_int))); build_opts.emplace(("-DB_VAL=" + support::cpp11::to_string(b_const_int))); - const int o1 = input->info()->quantization_info().offset; + const int o1 = input->info()->quantization_info().offset; + const float s1 = input->info()->quantization_info().scale; // Quantized value of 0 corresponds to the offset o1 build_opts.emplace(("-DCONST_0=" + support::cpp11::to_string(o1))); + build_opts.emplace(("-DS1_VAL=" + float_to_string_with_full_precision(s1))); + build_opts.emplace(("-DO1_VAL=" + support::cpp11::to_string(o1))); // Set scale and offset of the input and output if they have different quantization info if(is_data_type_quantized_asymmetric(dt) && output != nullptr) { - const float s1 = input->info()->quantization_info().scale; const float s2 = output->info()->quantization_info().scale; const int o2 = output->info()->quantization_info().offset; if(o1 != o2 || s1 != s2) { - build_opts.emplace(("-DS1_VAL=" + float_to_string_with_full_precision(s1))); build_opts.emplace(("-DS2_VAL=" + float_to_string_with_full_precision(s2))); - build_opts.emplace(("-DO1_VAL=" + support::cpp11::to_string(o1))); build_opts.emplace(("-DO2_VAL=" + support::cpp11::to_string(o2))); } } diff --git a/src/core/CL/kernels/CLDepthwiseConvolutionLayer3x3NCHWKernel.cpp b/src/core/CL/kernels/CLDepthwiseConvolutionLayer3x3NCHWKernel.cpp index 752a810520..8bd62c69f7 100644 --- a/src/core/CL/kernels/CLDepthwiseConvolutionLayer3x3NCHWKernel.cpp +++ b/src/core/CL/kernels/CLDepthwiseConvolutionLayer3x3NCHWKernel.cpp @@ -49,8 +49,9 @@ Status validate_arguments(const ITensorInfo *input, const ITensorInfo *weights, ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::QASYMM8, DataType::F16, DataType::F32); ARM_COMPUTE_RETURN_ERROR_ON_MSG(act_info.enabled() && ((input->data_type() != DataType::QASYMM8) || ((act_info.activation() != ActivationLayerInfo::ActivationFunction::LU_BOUNDED_RELU) && (act_info.activation() != ActivationLayerInfo::ActivationFunction::BOUNDED_RELU) - && (act_info.activation() != ActivationLayerInfo::ActivationFunction::RELU))), - "For QASYMM8 only relu, lower bounded relu and lower-upper bounded relu are supported"); + && (act_info.activation() != ActivationLayerInfo::ActivationFunction::RELU) + && (act_info.activation() != ActivationLayerInfo::ActivationFunction::LOGISTIC))), + "For QASYMM8 only logistic, relu, lower bounded relu and lower-upper bounded relu are supported"); ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_TYPES(input, weights); ARM_COMPUTE_RETURN_ERROR_ON(weights->dimension(0) != 3 || weights->dimension(1) != 3); ARM_COMPUTE_RETURN_ERROR_ON((input->dimension(2) * depth_multiplier) != output->dimension(2)); diff --git a/src/core/CL/kernels/CLDepthwiseConvolutionLayer3x3NHWCKernel.cpp b/src/core/CL/kernels/CLDepthwiseConvolutionLayer3x3NHWCKernel.cpp index 9d9c280182..7754e1b0e2 100644 --- a/src/core/CL/kernels/CLDepthwiseConvolutionLayer3x3NHWCKernel.cpp +++ b/src/core/CL/kernels/CLDepthwiseConvolutionLayer3x3NHWCKernel.cpp @@ -47,8 +47,9 @@ Status validate_arguments(const ITensorInfo *input, const ITensorInfo *weights, ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::F32, DataType::QASYMM8); ARM_COMPUTE_RETURN_ERROR_ON_MSG((act_info.enabled()) && (input->data_type() == DataType::F32 || ((act_info.activation() != ActivationLayerInfo::ActivationFunction::LU_BOUNDED_RELU) && (act_info.activation() != ActivationLayerInfo::ActivationFunction::BOUNDED_RELU) - && (act_info.activation() != ActivationLayerInfo::ActivationFunction::RELU))), - "For QASYMM8 only relu, lower bounded relu and lower-upper bounded relu are supported"); //COMPMID-1317 add fused activation for F32 + && (act_info.activation() != ActivationLayerInfo::ActivationFunction::RELU) + && (act_info.activation() != ActivationLayerInfo::ActivationFunction::LOGISTIC))), + "For QASYMM8 only logistic, relu, lower bounded relu and lower-upper bounded relu are supported"); //COMPMID-1317 add fused activation for F32 ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_TYPES(input, weights); ARM_COMPUTE_RETURN_ERROR_ON(depth_multiplier > 1); // COMPMID-1071 Add depth multiplier support for NHWC ARM_COMPUTE_RETURN_ERROR_ON(weights->dimension(1) != 3 || weights->dimension(2) != 3); -- cgit v1.2.1