From 8b2a7d3aa119e7f1d6a03690d05eb27c5d178b9f Mon Sep 17 00:00:00 2001 From: Giorgio Arena Date: Tue, 11 Feb 2020 17:21:31 +0000 Subject: COMPMID-3101 Fuse activation with floating point elementwise operation layers in CL Signed-off-by: Giorgio Arena Change-Id: I1693f8664ba7c0dc8c076bbe7365cef1e667bd25 Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/2718 Tested-by: Arm Jenkins Reviewed-by: Gian Marco Iodice Comments-Addressed: Arm Jenkins --- .../core/CL/kernels/CLElementwiseOperationKernel.h | 53 ++--- .../CL/kernels/CLPixelWiseMultiplicationKernel.h | 24 ++- arm_compute/graph/LayerDescriptors.h | 23 ++- arm_compute/graph/backends/FunctionHelpers.h | 7 +- arm_compute/graph/nodes/EltwiseLayerNode.h | 14 ++ .../runtime/CL/functions/CLElementwiseOperations.h | 164 ++++++++------- .../CL/functions/CLPixelWiseMultiplication.h | 28 +-- .../GLES_COMPUTE/functions/GCArithmeticAddition.h | 24 ++- .../functions/GCPixelWiseMultiplication.h | 13 +- .../runtime/NEON/functions/NEArithmeticAddition.h | 24 ++- .../NEON/functions/NEArithmeticSubtraction.h | 22 +- .../NEON/functions/NEElementwiseOperations.h | 90 ++++---- .../NEON/functions/NEPixelWiseMultiplication.h | 30 +-- src/core/CL/cl_kernels/activation_float_helpers.h | 4 +- src/core/CL/cl_kernels/elementwise_operation.cl | 12 +- src/core/CL/cl_kernels/pixelwise_mul_float.cl | 15 +- .../CL/kernels/CLElementwiseOperationKernel.cpp | 28 ++- .../CL/kernels/CLPixelWiseMultiplicationKernel.cpp | 38 +++- src/graph/mutators/NodeFusionMutator.cpp | 18 +- src/graph/nodes/EltwiseLayerNode.cpp | 10 + .../CL/functions/CLElementwiseOperations.cpp | 56 ++--- .../CL/functions/CLPixelWiseMultiplication.cpp | 16 +- .../functions/GCArithmeticAddition.cpp | 6 +- .../functions/GCPixelWiseMultiplication.cpp | 3 +- .../NEON/functions/NEArithmeticAddition.cpp | 6 +- .../NEON/functions/NEArithmeticSubtraction.cpp | 6 +- .../NEON/functions/NEElementwiseOperators.cpp | 30 ++- .../NEON/functions/NEPixelWiseMultiplication.cpp | 14 +- tests/datasets/ShapeDatasets.h | 23 ++- tests/validation/CL/ArithmeticAddition.cpp | 65 ++++-- tests/validation/CL/ArithmeticDivision.cpp | 51 ++++- tests/validation/CL/ArithmeticSubtraction.cpp | 65 ++++-- tests/validation/CL/ElementwiseMax.cpp | 55 +++-- tests/validation/CL/ElementwiseMin.cpp | 55 +++-- tests/validation/CL/ElementwisePower.cpp | 55 ++++- tests/validation/CL/ElementwiseSquaredDiff.cpp | 44 +++- tests/validation/CL/PixelWiseMultiplication.cpp | 32 ++- tests/validation/NEON/ElementwiseDivision.cpp | 6 +- .../fixtures/ArithmeticOperationsFixture.h | 72 ++++++- .../fixtures/ElementwiseOperationsFixture.h | 229 +++++++++++++++++---- .../fixtures/PixelWiseMultiplicationFixture.h | 66 ++++-- tests/validation/reference/ActivationLayer.cpp | 3 +- tests/validation/reference/ActivationLayer.h | 2 +- 43 files changed, 1127 insertions(+), 474 deletions(-) diff --git a/arm_compute/core/CL/kernels/CLElementwiseOperationKernel.h b/arm_compute/core/CL/kernels/CLElementwiseOperationKernel.h index 34fafaa3a5..85961f28bc 100644 --- a/arm_compute/core/CL/kernels/CLElementwiseOperationKernel.h +++ b/arm_compute/core/CL/kernels/CLElementwiseOperationKernel.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2019 ARM Limited. + * Copyright (c) 2018-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -97,6 +97,8 @@ protected: */ void configure_common(const ICLTensor *input1, const ICLTensor *input2, ICLTensor *output); + ActivationLayerInfo _act_info; + private: const ICLTensor *_input1; /**< Source tensor 1 */ const ICLTensor *_input2; /**< Source tensor 2 */ @@ -114,25 +116,28 @@ public: /** Static function to check if given info will lead to a valid configuration of @ref CLSaturatedArithmeticOperationKernel * - * @param[in] op Arithmetic operation to be executed. - * @param[in] input1 First tensor input. Data types supported: U8/S8/QASYMM8/QASYMM8_SIGNED/U16/S16/QSYMM16/F16/U32/S32/F32. - * @param[in] input2 Second tensor input. Data types supported: Same as @p input1. - * @param[in] output Output tensor. Data types supported: Same as @p input1. - * @param[in] policy Policy to use to handle overflow. + * @param[in] op Arithmetic operation to be executed. + * @param[in] input1 First tensor input. Data types supported: U8/S8/QASYMM8/QASYMM8_SIGNED/U16/S16/QSYMM16/F16/U32/S32/F32. + * @param[in] input2 Second tensor input. Data types supported: Same as @p input1. + * @param[in] output Output tensor. Data types supported: Same as @p input1. + * @param[in] policy Policy to use to handle overflow. + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. */ - void configure(ArithmeticOperation op, const ICLTensor *input1, const ICLTensor *input2, ICLTensor *output, const ConvertPolicy &policy); + void configure(ArithmeticOperation op, const ICLTensor *input1, const ICLTensor *input2, ICLTensor *output, const ConvertPolicy &policy, const ActivationLayerInfo &act_info = ActivationLayerInfo()); /** Static function to check if given info will lead to a valid configuration of @ref CLSaturatedArithmeticOperationKernel * - * @param[in] op Arithmetic operation to be executed. - * @param[in] input1 First tensor input info. Data types supported: U8/S8/QASYMM8/QASYMM8_SIGNED/U16/S16/QSYMM16/F16/U32/S32/F32. - * @param[in] input2 Second tensor input info. Data types supported: Same as @p input1. - * @param[in] output Output tensor info. Data types supported: Same as @p input1. - * @param[in] policy Policy to use to handle overflow. + * @param[in] op Arithmetic operation to be executed. + * @param[in] input1 First tensor input info. Data types supported: U8/S8/QASYMM8/QASYMM8_SIGNED/U16/S16/QSYMM16/F16/U32/S32/F32. + * @param[in] input2 Second tensor input info. Data types supported: Same as @p input1. + * @param[in] output Output tensor info. Data types supported: Same as @p input1. + * @param[in] policy Policy to use to handle overflow. + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. * * @return a Status */ - static Status validate(ArithmeticOperation op, const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, const ConvertPolicy &policy); + static Status validate(ArithmeticOperation op, const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, const ConvertPolicy &policy, + const ActivationLayerInfo &act_info = ActivationLayerInfo()); protected: // Inherited methods overridden: @@ -157,23 +162,25 @@ public: /** Static function to check if given info will lead to a valid configuration of @ref CLArithmeticOperationKernel * - * @param[in] op Arithmetic operation to be executed. - * @param[in] input1 First tensor input. Data types supported: U8/S8/QASYMM8/QASYMM8_SIGNED/U16/S16/QSYMM16/F16/U32/S32/F32. - * @param[in] input2 Second tensor input. Data types supported: Same as @p input1. - * @param[in] output Output tensor. Data types supported: Same as @p input1. + * @param[in] op Arithmetic operation to be executed. + * @param[in] input1 First tensor input. Data types supported: U8/S8/QASYMM8/QASYMM8_SIGNED/U16/S16/QSYMM16/F16/U32/S32/F32. + * @param[in] input2 Second tensor input. Data types supported: Same as @p input1. + * @param[in] output Output tensor. Data types supported: Same as @p input1. + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. */ - void configure(ArithmeticOperation op, const ICLTensor *input1, const ICLTensor *input2, ICLTensor *output); + void configure(ArithmeticOperation op, const ICLTensor *input1, const ICLTensor *input2, ICLTensor *output, const ActivationLayerInfo &act_info = ActivationLayerInfo()); /** Static function to check if given info will lead to a valid configuration of @ref CLArithmeticOperationKernel * - * @param[in] op Arithmetic operation to be executed. - * @param[in] input1 First tensor input info. Data types supported: U8/S8/QASYMM8/QASYMM8_SIGNED/U16/S16/QSYMM16/F16/U32/S32/F32. - * @param[in] input2 Second tensor input info. Data types supported: Same as @p input1. - * @param[in] output Output tensor info. Data types supported: Same as @p input1. + * @param[in] op Arithmetic operation to be executed. + * @param[in] input1 First tensor input info. Data types supported: U8/S8/QASYMM8/QASYMM8_SIGNED/U16/S16/QSYMM16/F16/U32/S32/F32. + * @param[in] input2 Second tensor input info. Data types supported: Same as @p input1. + * @param[in] output Output tensor info. Data types supported: Same as @p input1. + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. * * @return a Status */ - static Status validate(ArithmeticOperation op, const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output); + static Status validate(ArithmeticOperation op, const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, const ActivationLayerInfo &act_info = ActivationLayerInfo()); protected: // Inherited methods overridden: diff --git a/arm_compute/core/CL/kernels/CLPixelWiseMultiplicationKernel.h b/arm_compute/core/CL/kernels/CLPixelWiseMultiplicationKernel.h index 58471ab299..eacdb44c09 100644 --- a/arm_compute/core/CL/kernels/CLPixelWiseMultiplicationKernel.h +++ b/arm_compute/core/CL/kernels/CLPixelWiseMultiplicationKernel.h @@ -55,9 +55,10 @@ public: * Scale must be positive and its value must be either 1/255 or 1/2^n where n is between 0 and 15. * @param[in] overflow_policy Overflow policy. Supported overflow policies: Wrap, Saturate * @param[in] rounding_policy Rounding policy. Supported rounding modes: to zero, to nearest even. + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. */ void configure(const ICLTensor *input1, const ICLTensor *input2, ICLTensor *output, float scale, - ConvertPolicy overflow_policy, RoundingPolicy rounding_policy); + ConvertPolicy overflow_policy, RoundingPolicy rounding_policy, const ActivationLayerInfo &act_info = ActivationLayerInfo()); /** Static function to check if given info will lead to a valid configuration of @ref CLPixelWiseMultiplicationKernel * * @param[in] input1 An input tensor info. Data types supported: U8/QASYMM8/QASYMM8_SIGNED/S16/QSYMM16/F16/F32. @@ -67,11 +68,12 @@ public: * Scale must be positive and its value must be either 1/255 or 1/2^n where n is between 0 and 15. * @param[in] overflow_policy Overflow policy. Supported overflow policies: Wrap, Saturate * @param[in] rounding_policy Rounding policy. Supported rounding modes: to zero, to nearest even. + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. * * @return a status */ static Status validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, float scale, - ConvertPolicy overflow_policy, RoundingPolicy rounding_policy); + ConvertPolicy overflow_policy, RoundingPolicy rounding_policy, const ActivationLayerInfo &act_info = ActivationLayerInfo()); // Inherited methods overridden: void run(const Window &window, cl::CommandQueue &queue) override; @@ -99,20 +101,22 @@ public: CLComplexPixelWiseMultiplicationKernel &operator=(CLComplexPixelWiseMultiplicationKernel &&) = default; /** Initialise the kernel's input, output and border mode. * - * @param[in] input1 An input tensor. Data types supported: F32. Number of channels supported: 2. - * @param[in] input2 An input tensor. Data types supported: same as @p input1. Number of channels supported: same as @p input1. - * @param[out] output The output tensor, Data types supported: same as @p input1. Number of channels supported: same as @p input1. + * @param[in] input1 An input tensor. Data types supported: F32. Number of channels supported: 2. + * @param[in] input2 An input tensor. Data types supported: same as @p input1. Number of channels supported: same as @p input1. + * @param[out] output The output tensor, Data types supported: same as @p input1. Number of channels supported: same as @p input1. + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. */ - void configure(const ICLTensor *input1, const ICLTensor *input2, ICLTensor *output); + void configure(const ICLTensor *input1, const ICLTensor *input2, ICLTensor *output, const ActivationLayerInfo &act_info = ActivationLayerInfo()); /** Static function to check if given info will lead to a valid configuration of @ref CLComplexPixelWiseMultiplicationKernel * - * @param[in] input1 An input tensor info. Data types supported: F32. Number of channels supported: 2. - * @param[in] input2 An input tensor info. Data types supported: same as @p input1. Number of channels supported: same as @p input1. - * @param[in] output The output tensor info. Data types supported: same as @p input1. Number of channels supported: same as @p input1. + * @param[in] input1 An input tensor info. Data types supported: F32. Number of channels supported: 2. + * @param[in] input2 An input tensor info. Data types supported: same as @p input1. Number of channels supported: same as @p input1. + * @param[in] output The output tensor info. Data types supported: same as @p input1. Number of channels supported: same as @p input1. + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. * * @return a status */ - static Status validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output); + static Status validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, const ActivationLayerInfo &act_info = ActivationLayerInfo()); // Inherited methods overridden: void run(const Window &window, cl::CommandQueue &queue) override; diff --git a/arm_compute/graph/LayerDescriptors.h b/arm_compute/graph/LayerDescriptors.h index 0cf203174e..d8e6a6a87b 100644 --- a/arm_compute/graph/LayerDescriptors.h +++ b/arm_compute/graph/LayerDescriptors.h @@ -70,20 +70,23 @@ struct EltwiseLayerDescriptor { /** Constructor * - * @param[in] op Element-wise operation to perform - * @param[in] out_quant_info (Optional) Output quantization information. Defaults to empty @ref QuantizationInfo - * @param[in] c_policy (Optional) Convert policy used for the operation. Defaults to @ref ConvertPolicy::SATURATE - * @param[in] r_policy (Optional) Rounding policy used for the operation. Defaults to @ref RoundingPolicy::TO_ZERO + * @param[in] op Element-wise operation to perform + * @param[in] out_quant_info (Optional) Output quantization information. Defaults to empty @ref QuantizationInfo + * @param[in] c_policy (Optional) Convert policy used for the operation. Defaults to @ref ConvertPolicy::SATURATE + * @param[in] r_policy (Optional) Rounding policy used for the operation. Defaults to @ref RoundingPolicy::TO_ZERO + * @param[in] fused_activation (Optional) Fused activation information. Defaults to empty (identity) @ref ActivationLayerInfo */ - EltwiseLayerDescriptor(EltwiseOperation op, QuantizationInfo out_quant_info = QuantizationInfo(), ConvertPolicy c_policy = ConvertPolicy::SATURATE, RoundingPolicy r_policy = RoundingPolicy::TO_ZERO) - : op(op), out_quant_info(out_quant_info), c_policy(c_policy), r_policy(r_policy) + EltwiseLayerDescriptor(EltwiseOperation op, QuantizationInfo out_quant_info = QuantizationInfo(), ConvertPolicy c_policy = ConvertPolicy::SATURATE, RoundingPolicy r_policy = RoundingPolicy::TO_ZERO, + ActivationLayerInfo fused_activation = ActivationLayerInfo()) + : op(op), out_quant_info(out_quant_info), c_policy(c_policy), r_policy(r_policy), fused_activation(fused_activation) { } - EltwiseOperation op; /**< Element-wise operation to perform */ - QuantizationInfo out_quant_info; /**< Output quantization information */ - ConvertPolicy c_policy; /**< Convert policy */ - RoundingPolicy r_policy; /**< Rounding policy */ + EltwiseOperation op; /**< Element-wise operation to perform */ + QuantizationInfo out_quant_info; /**< Output quantization information */ + ConvertPolicy c_policy; /**< Convert policy */ + RoundingPolicy r_policy; /**< Rounding policy */ + ActivationLayerInfo fused_activation; /**< Fused activation info */ }; /** Deconvolution layer descriptor */ diff --git a/arm_compute/graph/backends/FunctionHelpers.h b/arm_compute/graph/backends/FunctionHelpers.h index 44b24b58bf..382b18a888 100644 --- a/arm_compute/graph/backends/FunctionHelpers.h +++ b/arm_compute/graph/backends/FunctionHelpers.h @@ -773,6 +773,7 @@ std::unique_ptr create_eltwise_layer(EltwiseLayerNode &node) typename TargetInfo::TensorType *output = get_backing_tensor(node.output(0)); const EltwiseOperation eltwise_op = node.eltwise_operation(); const ConvertPolicy convert_policy = node.convert_policy(); + const ActivationLayerInfo act_info = node.fused_activation(); ARM_COMPUTE_ERROR_ON(input1 == nullptr); ARM_COMPUTE_ERROR_ON(input2 == nullptr); ARM_COMPUTE_ERROR_ON(output == nullptr); @@ -783,19 +784,19 @@ std::unique_ptr create_eltwise_layer(EltwiseLayerNode &node) { std::tie(func, func_name) = create_named_function( std::string("ArithmeticAddition"), - input1, input2, output, convert_policy); + input1, input2, output, convert_policy, act_info); } else if(eltwise_op == EltwiseOperation::Sub) { std::tie(func, func_name) = create_named_function( std::string("ArithmeticSubtraction"), - input1, input2, output, convert_policy); + input1, input2, output, convert_policy, act_info); } else if(eltwise_op == EltwiseOperation::Mul) { std::tie(func, func_name) = create_named_function( std::string("PixelWiseMultiplication"), - input1, input2, output, 1.f, convert_policy, node.rounding_policy()); + input1, input2, output, 1.f, convert_policy, node.rounding_policy(), act_info); } else { diff --git a/arm_compute/graph/nodes/EltwiseLayerNode.h b/arm_compute/graph/nodes/EltwiseLayerNode.h index 21c220a548..d619ad2588 100644 --- a/arm_compute/graph/nodes/EltwiseLayerNode.h +++ b/arm_compute/graph/nodes/EltwiseLayerNode.h @@ -57,12 +57,26 @@ public: */ RoundingPolicy rounding_policy() const; + /** Returns fused activation + * + * @return Fused activation + */ + ActivationLayerInfo fused_activation() const; + + /** Sets fused activation + * + * @param[in] fused_activation Fused activation to set + */ + void set_fused_activation(ActivationLayerInfo fused_activation); + // Inherited overridden methods: NodeType type() const override; bool forward_descriptors() override; TensorDescriptor configure_output(size_t idx) const override; void accept(INodeVisitor &v) override; + static constexpr NodeType node_type = NodeType::EltwiseLayer; + private: descriptors::EltwiseLayerDescriptor descriptor; }; diff --git a/arm_compute/runtime/CL/functions/CLElementwiseOperations.h b/arm_compute/runtime/CL/functions/CLElementwiseOperations.h index a7cb8b4226..6d9f3a0e97 100644 --- a/arm_compute/runtime/CL/functions/CLElementwiseOperations.h +++ b/arm_compute/runtime/CL/functions/CLElementwiseOperations.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2019 ARM Limited. + * Copyright (c) 2018-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -41,24 +41,26 @@ class CLArithmeticAddition : public ICLSimpleFunction public: /** Initialise the kernel's inputs, output and conversion policy. * - * @param[in, out] input1 First tensor input. Data types supported: U8/QASYMM8/S16/QSYMM16/S32/U32/F16/F32. - * The input tensor is [in, out] because its TensorInfo might be modified inside the kernel in case of broadcasting of dimension 0. - * @param[in, out] input2 Second tensor input. Data types supported: U8, QASYMM8 (only if @p input1 is QASYMM8), QSYMM16 (only if @p input1 is QSYMM16), S16/F16/F32. - * The input tensor is [in, out] because its TensorInfo might be modified inside the kernel in case of broadcasting of dimension 0. - * @param[out] output Output tensor. Data types supported: U8 (Only if both inputs are U8), QASYMM8 (only if both inputs are QASYMM8), QSYMM16 (only if both inputs is QSYMM16), S16/F16/F32. - * @param[in] policy Policy to use to handle overflow. + * @param[in, out] input1 First tensor input. Data types supported: U8/QASYMM8/S16/QSYMM16/S32/U32/F16/F32. + * The input tensor is [in, out] because its TensorInfo might be modified inside the kernel in case of broadcasting of dimension 0. + * @param[in, out] input2 Second tensor input. Data types supported: U8, QASYMM8 (only if @p input1 is QASYMM8), QSYMM16 (only if @p input1 is QSYMM16), S16/F16/F32. + * The input tensor is [in, out] because its TensorInfo might be modified inside the kernel in case of broadcasting of dimension 0. + * @param[out] output Output tensor. Data types supported: U8 (Only if both inputs are U8), QASYMM8 (only if both inputs are QASYMM8), QSYMM16 (only if both inputs is QSYMM16), S16/F16/F32. + * @param[in] policy Policy to use to handle overflow. + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. */ - void configure(ICLTensor *input1, ICLTensor *input2, ICLTensor *output, ConvertPolicy policy); + void configure(ICLTensor *input1, ICLTensor *input2, ICLTensor *output, ConvertPolicy policy, const ActivationLayerInfo &act_info = ActivationLayerInfo()); /** Static function to check if given info will lead to a valid configuration of @ref CLSaturatedArithmeticOperationKernel for addition * - * @param[in] input1 First tensor input info. Data types supported: U8/QASYMM8/S16/QSYMM16/S32/U32/F16/F32. - * @param[in] input2 Second tensor input info. Data types supported: U8, QASYMM8 (only if @p input1 is QASYMM8), QSYMM16 (only if @p input1 is QSYMM16), S16/F16/F32. - * @param[in] output Output tensor info. Data types supported: U8 (Only if both inputs are U8), QASYMM8 (only if both inputs are QASYMM8), QSYMM16 (only if both inputs is QSYMM16), S16/F16/F32. - * @param[in] policy Policy to use to handle overflow. + * @param[in] input1 First tensor input info. Data types supported: U8/QASYMM8/S16/QSYMM16/S32/U32/F16/F32. + * @param[in] input2 Second tensor input info. Data types supported: U8, QASYMM8 (only if @p input1 is QASYMM8), QSYMM16 (only if @p input1 is QSYMM16), S16/F16/F32. + * @param[in] output Output tensor info. Data types supported: U8 (Only if both inputs are U8), QASYMM8 (only if both inputs are QASYMM8), QSYMM16 (only if both inputs is QSYMM16), S16/F16/F32. + * @param[in] policy Policy to use to handle overflow. + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. * * @return a status */ - static Status validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, ConvertPolicy policy); + static Status validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, ConvertPolicy policy, const ActivationLayerInfo &act_info = ActivationLayerInfo()); }; /** Basic function to run @ref CLSaturatedArithmeticOperationKernel for subtraction @@ -71,24 +73,26 @@ class CLArithmeticSubtraction : public ICLSimpleFunction public: /** Initialise the kernel's inputs, output and conversion policy. * - * @param[in, out] input1 First tensor input. Data types supported: U8/QASYMM8/S16/S32/U32/F16/F32. - * The input tensor is [in, out] because its TensorInfo might be modified inside the kernel in case of broadcasting of dimension 0. - * @param[in, out] input2 Second tensor input. Data types supported: U8, QASYMM8 (only if @p input1 is QASYMM8), S16/F16/F32. - * The input tensor is [in, out] because its TensorInfo might be modified inside the kernel in case of broadcasting of dimension 0. - * @param[out] output Output tensor. Data types supported: U8 (Only if both inputs are U8), QASYMM8 (only if both inputs are QASYMM8), S16/F16/F32. - * @param[in] policy Policy to use to handle overflow. + * @param[in, out] input1 First tensor input. Data types supported: U8/QASYMM8/S16/S32/U32/F16/F32. + * The input tensor is [in, out] because its TensorInfo might be modified inside the kernel in case of broadcasting of dimension 0. + * @param[in, out] input2 Second tensor input. Data types supported: U8, QASYMM8 (only if @p input1 is QASYMM8), S16/F16/F32. + * The input tensor is [in, out] because its TensorInfo might be modified inside the kernel in case of broadcasting of dimension 0. + * @param[out] output Output tensor. Data types supported: U8 (Only if both inputs are U8), QASYMM8 (only if both inputs are QASYMM8), S16/F16/F32. + * @param[in] policy Policy to use to handle overflow. + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. */ - void configure(ICLTensor *input1, ICLTensor *input2, ICLTensor *output, ConvertPolicy policy); + void configure(ICLTensor *input1, ICLTensor *input2, ICLTensor *output, ConvertPolicy policy, const ActivationLayerInfo &act_info = ActivationLayerInfo()); /** Static function to check if given info will lead to a valid configuration of @ref CLSaturatedArithmeticOperationKernel for subtraction * - * @param[in] input1 First tensor input info. Data types supported: U8/QASYMM8/S16/S32/U32/F16/F32. - * @param[in] input2 Second tensor input info. Data types supported: U8, QASYMM8 (only if @p input1 is QASYMM8), S16/F16/F32. - * @param[in] output Output tensor info. Data types supported: U8 (Only if both inputs are U8), QASYMM8 ( only if both inputs are QASYMM8), S16/F16/F32. - * @param[in] policy Policy to use to handle overflow. + * @param[in] input1 First tensor input info. Data types supported: U8/QASYMM8/S16/S32/U32/F16/F32. + * @param[in] input2 Second tensor input info. Data types supported: U8, QASYMM8 (only if @p input1 is QASYMM8), S16/F16/F32. + * @param[in] output Output tensor info. Data types supported: U8 (Only if both inputs are U8), QASYMM8 ( only if both inputs are QASYMM8), S16/F16/F32. + * @param[in] policy Policy to use to handle overflow. + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. * * @return a status */ - static Status validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, ConvertPolicy policy); + static Status validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, ConvertPolicy policy, const ActivationLayerInfo &act_info = ActivationLayerInfo()); }; /** Basic function to run @ref CLSaturatedArithmeticOperationKernel for division @@ -101,22 +105,24 @@ class CLArithmeticDivision : public ICLSimpleFunction public: /** Initialise the kernel's inputs, output. * - * @param[in, out] input1 First tensor input. Data types supported: F16/F32. - * The input tensor is [in, out] because its TensorInfo might be modified inside the kernel in case of broadcasting of dimension 0. - * @param[in, out] input2 Second tensor input. Same as @p input1. - * The input tensor is [in, out] because its TensorInfo might be modified inside the kernel in case of broadcasting of dimension 0. - * @param[out] output Output tensor. Data types supported: Same as @p input1. + * @param[in, out] input1 First tensor input. Data types supported: F16/F32. + * The input tensor is [in, out] because its TensorInfo might be modified inside the kernel in case of broadcasting of dimension 0. + * @param[in, out] input2 Second tensor input. Same as @p input1. + * The input tensor is [in, out] because its TensorInfo might be modified inside the kernel in case of broadcasting of dimension 0. + * @param[out] output Output tensor. Data types supported: Same as @p input1. + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. */ - void configure(ICLTensor *input1, ICLTensor *input2, ICLTensor *output); + void configure(ICLTensor *input1, ICLTensor *input2, ICLTensor *output, const ActivationLayerInfo &act_info = ActivationLayerInfo()); /** Static function to check if given info will lead to a valid configuration of @ref CLArithmeticDivision * - * @param[in] input1 First tensor input info. Data types supported: F16/F32. - * @param[in] input2 Second tensor input info. Data types supported: Same as @p input1. - * @param[in] output Output tensor info. Data types supported: Same as @p input1. + * @param[in] input1 First tensor input info. Data types supported: F16/F32. + * @param[in] input2 Second tensor input info. Data types supported: Same as @p input1. + * @param[in] output Output tensor info. Data types supported: Same as @p input1. + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. * * @return a status */ - static Status validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output); + static Status validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, const ActivationLayerInfo &act_info = ActivationLayerInfo()); }; /** Basic function to run @ref CLArithmeticOperationKernel for max @@ -129,22 +135,24 @@ class CLElementwiseMax : public ICLSimpleFunction public: /** Initialise the kernel's inputs, output and conversion policy. * - * @param[in, out] input1 First tensor input. Data types supported: U8/QASYMM8/S16/QSYMM16/S32/U32/F16/F32. - * The input tensor is [in, out] because its TensorInfo might be modified inside the kernel in case of broadcasting of dimension 0. - * @param[in, out] input2 Second tensor input. Data types supported: U8, QASYMM8 (only if @p input1 is QASYMM8), S16, QSYMM16 (only if @p input1 is QSYMM16), F16/F32. - * The input tensor is [in, out] because its TensorInfo might be modified inside the kernel in case of broadcasting of dimension 0. - * @param[out] output Output tensor. Data types supported: U8 (Only if both inputs are U8), QASYMM8 (only if both inputs are QASYMM8), S16, QSYMM16 (only if both inputs are QSYMM16), F16/F32. + * @param[in, out] input1 First tensor input. Data types supported: U8/QASYMM8/S16/QSYMM16/S32/U32/F16/F32. + * The input tensor is [in, out] because its TensorInfo might be modified inside the kernel in case of broadcasting of dimension 0. + * @param[in, out] input2 Second tensor input. Data types supported: U8, QASYMM8 (only if @p input1 is QASYMM8), S16, QSYMM16 (only if @p input1 is QSYMM16), F16/F32. + * The input tensor is [in, out] because its TensorInfo might be modified inside the kernel in case of broadcasting of dimension 0. + * @param[out] output Output tensor. Data types supported: U8 (Only if both inputs are U8), QASYMM8 (only if both inputs are QASYMM8), S16, QSYMM16 (only if both inputs are QSYMM16), F16/F32. + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. */ - void configure(ICLTensor *input1, ICLTensor *input2, ICLTensor *output); + void configure(ICLTensor *input1, ICLTensor *input2, ICLTensor *output, const ActivationLayerInfo &act_info = ActivationLayerInfo()); /** Static function to check if given info will lead to a valid configuration of @ref CLArithmeticOperationKernel for max * - * @param[in] input1 First tensor input info. Data types supported: U8/QASYMM8/S16/QSYMM16/S32/U32/F16/F32. - * @param[in] input2 Second tensor input info. Data types supported: U8, QASYMM8 (only if @p input1 is QASYMM8), S16, QSYMM16 (only if @p input1 is QSYMM16), F16/F32. - * @param[in] output Output tensor info. Data types supported: U8 (Only if both inputs are U8), QASYMM8 ( only if both inputs are QASYMM8), S16, QSYMM16 (only if both inputs are QSYMM16), F16/F32. + * @param[in] input1 First tensor input info. Data types supported: U8/QASYMM8/S16/QSYMM16/S32/U32/F16/F32. + * @param[in] input2 Second tensor input info. Data types supported: U8, QASYMM8 (only if @p input1 is QASYMM8), S16, QSYMM16 (only if @p input1 is QSYMM16), F16/F32. + * @param[in] output Output tensor info. Data types supported: U8 (Only if both inputs are U8), QASYMM8 ( only if both inputs are QASYMM8), S16, QSYMM16 (only if both inputs are QSYMM16), F16/F32. + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. * * @return a status */ - static Status validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output); + static Status validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, const ActivationLayerInfo &act_info = ActivationLayerInfo()); }; /** Basic function to run @ref CLArithmeticOperationKernel for min @@ -157,22 +165,24 @@ class CLElementwiseMin : public ICLSimpleFunction public: /** Initialise the kernel's inputs, output and conversion policy. * - * @param[in, out] input1 First tensor input. Data types supported: U8/QASYMM8/S16/QSYMM16/S32/U32/F16/F32. - * The input tensor is [in, out] because its TensorInfo might be modified inside the kernel in case of broadcasting of dimension 0. - * @param[in, out] input2 Second tensor input. Data types supported: U8, QASYMM8 (only if @p input1 is QASYMM8), S16, QSYMM16 (only if @p input1 is QSYMM16), F16/F32. - * The input tensor is [in, out] because its TensorInfo might be modified inside the kernel in case of broadcasting of dimension 0. - * @param[out] output Output tensor. Data types supported: U8 (Only if both inputs are U8), QASYMM8 (only if both inputs are QASYMM8), S16, QSYMM16 (only if both inputs are QSYMM16), F16/F32. + * @param[in, out] input1 First tensor input. Data types supported: U8/QASYMM8/S16/QSYMM16/S32/U32/F16/F32. + * The input tensor is [in, out] because its TensorInfo might be modified inside the kernel in case of broadcasting of dimension 0. + * @param[in, out] input2 Second tensor input. Data types supported: U8, QASYMM8 (only if @p input1 is QASYMM8), S16, QSYMM16 (only if @p input1 is QSYMM16), F16/F32. + * The input tensor is [in, out] because its TensorInfo might be modified inside the kernel in case of broadcasting of dimension 0. + * @param[out] output Output tensor. Data types supported: U8 (Only if both inputs are U8), QASYMM8 (only if both inputs are QASYMM8), S16, QSYMM16 (only if both inputs are QSYMM16), F16/F32. + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. */ - void configure(ICLTensor *input1, ICLTensor *input2, ICLTensor *output); + void configure(ICLTensor *input1, ICLTensor *input2, ICLTensor *output, const ActivationLayerInfo &act_info = ActivationLayerInfo()); /** Static function to check if given info will lead to a valid configuration of @ref CLArithmeticOperationKernel for min * - * @param[in] input1 First tensor input info. Data types supported: U8/QASYMM8/S16/QSYMM16/S32/U32/F16/F32. - * @param[in] input2 Second tensor input info. Data types supported: U8, QASYMM8 (only if @p input1 is QASYMM8), S16, QSYMM16 (only if @p input1 is QSYMM16), F16/F32. - * @param[in] output Output tensor info. Data types supported: U8 (Only if both inputs are U8), QASYMM8 ( only if both inputs are QASYMM8), S16, QSYMM16 (only if both inputs are QSYMM16), F16/F32. + * @param[in] input1 First tensor input info. Data types supported: U8/QASYMM8/S16/QSYMM16/S32/U32/F16/F32. + * @param[in] input2 Second tensor input info. Data types supported: U8, QASYMM8 (only if @p input1 is QASYMM8), S16, QSYMM16 (only if @p input1 is QSYMM16), F16/F32. + * @param[in] output Output tensor info. Data types supported: U8 (Only if both inputs are U8), QASYMM8 ( only if both inputs are QASYMM8), S16, QSYMM16 (only if both inputs are QSYMM16), F16/F32. + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. * * @return a status */ - static Status validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output); + static Status validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, const ActivationLayerInfo &act_info = ActivationLayerInfo()); }; /** Basic function to run @ref CLArithmeticOperationKernel for squared difference @@ -185,22 +195,24 @@ class CLElementwiseSquaredDiff : public ICLSimpleFunction public: /** Initialise the kernel's inputs, output and conversion policy. * - * @param[in, out] input1 First tensor input. Data types supported: U8/QASYMM8/S16/QSYMM16/F16/F32. - * The input tensor is [in, out] because its TensorInfo might be modified inside the kernel in case of broadcasting of dimension 0. - * @param[in, out] input2 Second tensor input. Data types supported: U8, QASYMM8 (only if @p input1 is QASYMM8), S16, QSYMM16 (only if @p input1 is QSYMM16), F16/F32. - * The input tensor is [in, out] because its TensorInfo might be modified inside the kernel in case of broadcasting of dimension 0. - * @param[out] output Output tensor. Data types supported: U8 (Only if both inputs are U8), QASYMM8 (only if both inputs are QASYMM8), S16, QSYMM16 (only if both inputs are QSYMM16), F16/F32. + * @param[in, out] input1 First tensor input. Data types supported: U8/QASYMM8/S16/QSYMM16/F16/F32. + * The input tensor is [in, out] because its TensorInfo might be modified inside the kernel in case of broadcasting of dimension 0. + * @param[in, out] input2 Second tensor input. Data types supported: U8, QASYMM8 (only if @p input1 is QASYMM8), S16, QSYMM16 (only if @p input1 is QSYMM16), F16/F32. + * The input tensor is [in, out] because its TensorInfo might be modified inside the kernel in case of broadcasting of dimension 0. + * @param[out] output Output tensor. Data types supported: U8 (Only if both inputs are U8), QASYMM8 (only if both inputs are QASYMM8), S16, QSYMM16 (only if both inputs are QSYMM16), F16/F32. + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. */ - void configure(ICLTensor *input1, ICLTensor *input2, ICLTensor *output); + void configure(ICLTensor *input1, ICLTensor *input2, ICLTensor *output, const ActivationLayerInfo &act_info = ActivationLayerInfo()); /** Static function to check if given info will lead to a valid configuration of @ref CLArithmeticOperationKernel for squared difference * - * @param[in] input1 First tensor input info. Data types supported: U8/QASYMM8/S16/QSYMM16/F16/F32. - * @param[in] input2 Second tensor input info. Data types supported: U8, QASYMM8 (only if @p input1 is QASYMM8), S16, QSYMM16 (only if @p input1 is QSYMM16), F16/F32. - * @param[in] output Output tensor info. Data types supported: U8 (Only if both inputs are U8), QASYMM8 ( only if both inputs are QASYMM8), S16, QSYMM16 (only if both inputs are QSYMM16), F16/F32. + * @param[in] input1 First tensor input info. Data types supported: U8/QASYMM8/S16/QSYMM16/F16/F32. + * @param[in] input2 Second tensor input info. Data types supported: U8, QASYMM8 (only if @p input1 is QASYMM8), S16, QSYMM16 (only if @p input1 is QSYMM16), F16/F32. + * @param[in] output Output tensor info. Data types supported: U8 (Only if both inputs are U8), QASYMM8 ( only if both inputs are QASYMM8), S16, QSYMM16 (only if both inputs are QSYMM16), F16/F32. + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. * * @return a status */ - static Status validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output); + static Status validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, const ActivationLayerInfo &act_info = ActivationLayerInfo()); }; /** Basic function to run @ref CLArithmeticOperationKernel for power @@ -213,22 +225,24 @@ class CLElementwisePower : public ICLSimpleFunction public: /** Initialise the kernel's inputs, output and conversion policy. * - * @param[in, out] input1 First tensor input. Data types supported: F16/F32. - * The input tensor is [in, out] because its TensorInfo might be modified inside the kernel in case of broadcasting of dimension 0. - * @param[in, out] input2 Second tensor input. Data types supported: F16/F32. - * The input tensor is [in, out] because its TensorInfo might be modified inside the kernel in case of broadcasting of dimension 0. - * @param[out] output Output tensor. Data types supported:F16/F32. + * @param[in, out] input1 First tensor input. Data types supported: F16/F32. + * The input tensor is [in, out] because its TensorInfo might be modified inside the kernel in case of broadcasting of dimension 0. + * @param[in, out] input2 Second tensor input. Data types supported: F16/F32. + * The input tensor is [in, out] because its TensorInfo might be modified inside the kernel in case of broadcasting of dimension 0. + * @param[out] output Output tensor. Data types supported:F16/F32. + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. */ - void configure(ICLTensor *input1, ICLTensor *input2, ICLTensor *output); + void configure(ICLTensor *input1, ICLTensor *input2, ICLTensor *output, const ActivationLayerInfo &act_info = ActivationLayerInfo()); /** Static function to check if given info will lead to a valid configuration of @ref CLArithmeticOperationKernel for power * - * @param[in] input1 First tensor input info. Data types supported: F16/F32. - * @param[in] input2 Second tensor input info. Data types supported: F16/F32. - * @param[in] output Output tensor info. Data types supported: F16/F32. + * @param[in] input1 First tensor input info. Data types supported: F16/F32. + * @param[in] input2 Second tensor input info. Data types supported: F16/F32. + * @param[in] output Output tensor info. Data types supported: F16/F32. + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. * * @return a status */ - static Status validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output); + static Status validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, const ActivationLayerInfo &act_info = ActivationLayerInfo()); }; } // namespace arm_compute #endif /* ARM_COMPUTE_CLELEMENTWISEOPERATIONS_H */ diff --git a/arm_compute/runtime/CL/functions/CLPixelWiseMultiplication.h b/arm_compute/runtime/CL/functions/CLPixelWiseMultiplication.h index 72b1587b02..a5ab829c83 100644 --- a/arm_compute/runtime/CL/functions/CLPixelWiseMultiplication.h +++ b/arm_compute/runtime/CL/functions/CLPixelWiseMultiplication.h @@ -47,9 +47,10 @@ public: * Scale must be positive and its value must be either 1/255 or 1/2^n where n is between 0 and 15. * @param[in] overflow_policy Overflow policy. Supported overflow policies: Wrap, Saturate * @param[in] rounding_policy Rounding policy. Supported rounding modes: to zero, to nearest even. + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. */ void configure(ICLTensor *input1, ICLTensor *input2, ICLTensor *output, float scale, - ConvertPolicy overflow_policy, RoundingPolicy rounding_policy); + ConvertPolicy overflow_policy, RoundingPolicy rounding_policy, const ActivationLayerInfo &act_info = ActivationLayerInfo()); /** Static function to check if given info will lead to a valid configuration of @ref CLPixelWiseMultiplication * * @param[in] input1 An input tensor info. Data types supported: U8/QASYMM8/QASYMM8_SIGNED/S16/QSYMM16/F16/F32. @@ -59,11 +60,12 @@ public: * Scale must be positive and its value must be either 1/255 or 1/2^n where n is between 0 and 15. * @param[in] overflow_policy Overflow policy. Supported overflow policies: Wrap, Saturate * @param[in] rounding_policy Rounding policy. Supported rounding modes: to zero, to nearest even. + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. * * @return a status */ static Status validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, float scale, - ConvertPolicy overflow_policy, RoundingPolicy rounding_policy); + ConvertPolicy overflow_policy, RoundingPolicy rounding_policy, const ActivationLayerInfo &act_info = ActivationLayerInfo()); }; /** Basic function to run @ref CLComplexPixelWiseMultiplicationKernel. */ @@ -72,20 +74,22 @@ class CLComplexPixelWiseMultiplication : public ICLSimpleFunction public: /** Initialise the kernel's inputs, output. * - * @param[in, out] input1 An input tensor. Data types supported: F32. Number of channels supported: 2. - * The input tensor is [in, out] because its TensorInfo might be modified inside the kernel in case of broadcasting of dimension 0. - * @param[in, out] input2 An input tensor. Data types supported: same as @p input1. Number of channels supported: same as @p input1. - * The input tensor is [in, out] because its TensorInfo might be modified inside the kernel in case of broadcasting of dimension 0. - * @param[out] output The output tensor, Data types supported: same as @p input1. Number of channels supported: same as @p input1. + * @param[in, out] input1 An input tensor. Data types supported: F32. Number of channels supported: 2. + * The input tensor is [in, out] because its TensorInfo might be modified inside the kernel in case of broadcasting of dimension 0. + * @param[in, out] input2 An input tensor. Data types supported: same as @p input1. Number of channels supported: same as @p input1. + * The input tensor is [in, out] because its TensorInfo might be modified inside the kernel in case of broadcasting of dimension 0. + * @param[out] output The output tensor, Data types supported: same as @p input1. Number of channels supported: same as @p input1. + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. */ - void configure(ICLTensor *input1, ICLTensor *input2, ICLTensor *output); + void configure(ICLTensor *input1, ICLTensor *input2, ICLTensor *output, const ActivationLayerInfo &act_info = ActivationLayerInfo()); /** Static function to check if given info will lead to a valid configuration of @ref CLComplexPixelWiseMultiplication * - * @param[in] input1 An input tensor info. Data types supported: F32. Number of channels supported: 2. - * @param[in] input2 An input tensor info. Data types supported: same as @p input1. Number of channels supported: same as @p input1. - * @param[in] output The output tensor info, Data types supported: same as @p input1. Number of channels supported: same as @p input1. + * @param[in] input1 An input tensor info. Data types supported: F32. Number of channels supported: 2. + * @param[in] input2 An input tensor info. Data types supported: same as @p input1. Number of channels supported: same as @p input1. + * @param[in] output The output tensor info, Data types supported: same as @p input1. Number of channels supported: same as @p input1. + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. */ - static Status validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output); + static Status validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, const ActivationLayerInfo &act_info = ActivationLayerInfo()); }; } // namespace arm_compute #endif /*ARM_COMPUTE_CLPIXELWISEMULTIPLICATION_H */ diff --git a/arm_compute/runtime/GLES_COMPUTE/functions/GCArithmeticAddition.h b/arm_compute/runtime/GLES_COMPUTE/functions/GCArithmeticAddition.h index a16ab2d1ab..65bbacf272 100644 --- a/arm_compute/runtime/GLES_COMPUTE/functions/GCArithmeticAddition.h +++ b/arm_compute/runtime/GLES_COMPUTE/functions/GCArithmeticAddition.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016-2019 ARM Limited. + * Copyright (c) 2016-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -41,22 +41,24 @@ class GCArithmeticAddition : public IGCSimpleFunction public: /** Initialise the kernel's inputs, output and convertion policy. * - * @param[in] input1 First tensor input. Data types supported: F16. - * @param[in] input2 Second tensor input. Data types supported: F16. - * @param[out] output Output tensor. Data types supported: F16. - * @param[in] policy Policy to use to handle overflow. + * @param[in] input1 First tensor input. Data types supported: F16. + * @param[in] input2 Second tensor input. Data types supported: F16. + * @param[out] output Output tensor. Data types supported: F16. + * @param[in] policy Policy to use to handle overflow. + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. Currently not supported. */ - void configure(const IGCTensor *input1, const IGCTensor *input2, IGCTensor *output, ConvertPolicy policy); + void configure(const IGCTensor *input1, const IGCTensor *input2, IGCTensor *output, ConvertPolicy policy, const ActivationLayerInfo &act_info = ActivationLayerInfo()); /** Static function to check if given info will lead to a valid configuration of @ref GCArithmeticAddition * - * @param[in] input1 First tensor input info. Data types supported: F16. - * @param[in] input2 Second tensor input info. Data types supported: F16. - * @param[in] output Output tensor info. Data types supported: F16. - * @param[in] policy Policy to use to handle overflow. + * @param[in] input1 First tensor input info. Data types supported: F16. + * @param[in] input2 Second tensor input info. Data types supported: F16. + * @param[in] output Output tensor info. Data types supported: F16. + * @param[in] policy Policy to use to handle overflow. + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. Currently not supported. * * @return a status */ - static Status validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, ConvertPolicy policy); + static Status validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, ConvertPolicy policy, const ActivationLayerInfo &act_info = ActivationLayerInfo()); }; } #endif /* ARM_COMPUTE_GCARITHMETICADDITION_H */ diff --git a/arm_compute/runtime/GLES_COMPUTE/functions/GCPixelWiseMultiplication.h b/arm_compute/runtime/GLES_COMPUTE/functions/GCPixelWiseMultiplication.h index 6baa0de501..201e131dd9 100644 --- a/arm_compute/runtime/GLES_COMPUTE/functions/GCPixelWiseMultiplication.h +++ b/arm_compute/runtime/GLES_COMPUTE/functions/GCPixelWiseMultiplication.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2019 ARM Limited. + * Copyright (c) 2017-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -37,12 +37,13 @@ class GCPixelWiseMultiplication : public IGCSimpleFunction public: /** Initialise the kernel's inputs, output and convertion policy. * - * @param[in] input1 First tensor input. Data types supported: F32. - * @param[in] input2 Second tensor input. Data types supported: Same as @p input1. - * @param[out] output Output tensor. Data types supported: Same as @p input1. - * @param[in] scale Scale to apply after multiplication. Must be a positive value. + * @param[in] input1 First tensor input. Data types supported: F32. + * @param[in] input2 Second tensor input. Data types supported: Same as @p input1. + * @param[out] output Output tensor. Data types supported: Same as @p input1. + * @param[in] scale Scale to apply after multiplication. Must be a positive value. + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. Currently not supported. */ - void configure(const IGCTensor *input1, const IGCTensor *input2, IGCTensor *output, float scale); + void configure(const IGCTensor *input1, const IGCTensor *input2, IGCTensor *output, float scale, const ActivationLayerInfo &act_info = ActivationLayerInfo()); }; } #endif /*ARM_COMPUTE_GCPIXELWISEMULTIPLICATION_H */ diff --git a/arm_compute/runtime/NEON/functions/NEArithmeticAddition.h b/arm_compute/runtime/NEON/functions/NEArithmeticAddition.h index e17c255d2a..6cab5b3547 100644 --- a/arm_compute/runtime/NEON/functions/NEArithmeticAddition.h +++ b/arm_compute/runtime/NEON/functions/NEArithmeticAddition.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016-2019 ARM Limited. + * Copyright (c) 2016-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -37,22 +37,24 @@ class NEArithmeticAddition : public INESimpleFunction public: /** Initialise the kernel's inputs, output and conversion policy. * - * @param[in] input1 First tensor input. Data types supported: U8/QASYMM8/QASYMM8_SIGNED/S16/QSYMM16/F16/F32 - * @param[in] input2 Second tensor input. Data types supported: U8/QASYMM8/QASYMM8_SIGNED/S16/QSYMM16/F16/F32 - * @param[out] output Output tensor. Data types supported: U8/QASYMM8/QASYMM8_SIGNED/S16/QSYMM16/F16/F32 - * @param[in] policy Policy to use to handle overflow. + * @param[in] input1 First tensor input. Data types supported: U8/QASYMM8/QASYMM8_SIGNED/S16/QSYMM16/F16/F32 + * @param[in] input2 Second tensor input. Data types supported: U8/QASYMM8/QASYMM8_SIGNED/S16/QSYMM16/F16/F32 + * @param[out] output Output tensor. Data types supported: U8/QASYMM8/QASYMM8_SIGNED/S16/QSYMM16/F16/F32 + * @param[in] policy Policy to use to handle overflow. + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. Currently not supported. */ - void configure(ITensor *input1, ITensor *input2, ITensor *output, ConvertPolicy policy); + void configure(ITensor *input1, ITensor *input2, ITensor *output, ConvertPolicy policy, const ActivationLayerInfo &act_info = ActivationLayerInfo()); /** Static function to check if given info will lead to a valid configuration of @ref NEArithmeticAddition * - * @param[in] input1 First tensor input. Data types supported: U8/QASYMM8/QASYMM8_SIGNED/S16/QSYMM16/F16/F32 - * @param[in] input2 Second tensor input. Data types supported: U8/QASYMM8/QASYMM8_SIGNED/S16/QSYMM16/F16/F32 - * @param[in] output Output tensor. Data types supported: U8/SQASYMM8/QASYMM8_SIGNED/S16/QSYMM16/F16/F32 - * @param[in] policy Policy to use to handle overflow. + * @param[in] input1 First tensor input. Data types supported: U8/QASYMM8/QASYMM8_SIGNED/S16/QSYMM16/F16/F32 + * @param[in] input2 Second tensor input. Data types supported: U8/QASYMM8/QASYMM8_SIGNED/S16/QSYMM16/F16/F32 + * @param[in] output Output tensor. Data types supported: U8/SQASYMM8/QASYMM8_SIGNED/S16/QSYMM16/F16/F32 + * @param[in] policy Policy to use to handle overflow + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. Currently not supported. * * @return a status */ - static Status validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, ConvertPolicy policy); + static Status validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, ConvertPolicy policy, const ActivationLayerInfo &act_info = ActivationLayerInfo()); }; } // namespace arm_compute #endif /*ARM_COMPUTE_NEARITHMETICADDITION_H */ diff --git a/arm_compute/runtime/NEON/functions/NEArithmeticSubtraction.h b/arm_compute/runtime/NEON/functions/NEArithmeticSubtraction.h index c8c3fd3d2f..69d7b4bcfb 100644 --- a/arm_compute/runtime/NEON/functions/NEArithmeticSubtraction.h +++ b/arm_compute/runtime/NEON/functions/NEArithmeticSubtraction.h @@ -45,22 +45,24 @@ class NEArithmeticSubtraction : public INESimpleFunction public: /** Initialise the kernel's inputs, output and conversion policy. * - * @param[in] input1 First tensor input. Data types supported: U8/QASYMM8/QASYMM8_SIGNED/QSYMM16/S16/F16/F32 - * @param[in] input2 Second tensor input. Data types supported: U8/QASYMM8/QASYMM8_SIGNED/QSYMM16/S16/F16/F32 - * @param[out] output Output tensor. Data types supported: U8/QASYMM8/QASYMM8_SIGNED/QSYMM16/S16/F16/F32 - * @param[in] policy Policy to use to handle overflow. Convert policy cannot be WRAP if datatype is quantized. + * @param[in] input1 First tensor input. Data types supported: U8/QASYMM8/QASYMM8_SIGNED/QSYMM16/S16/F16/F32 + * @param[in] input2 Second tensor input. Data types supported: U8/QASYMM8/QASYMM8_SIGNED/QSYMM16/S16/F16/F32 + * @param[out] output Output tensor. Data types supported: U8/QASYMM8/QASYMM8_SIGNED/QSYMM16/S16/F16/F32 + * @param[in] policy Policy to use to handle overflow. Convert policy cannot be WRAP if datatype is quantized. + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. Currently not supported. */ - void configure(ITensor *input1, ITensor *input2, ITensor *output, ConvertPolicy policy); + void configure(ITensor *input1, ITensor *input2, ITensor *output, ConvertPolicy policy, const ActivationLayerInfo &act_info = ActivationLayerInfo()); /** Static function to check if given info will lead to a valid configuration of @ref NEArithmeticSubtraction * - * @param[in] input1 First tensor input. Data types supported: U8/QASYMM8/QASYMM8_SIGNED/S16/F16/F32 - * @param[in] input2 Second tensor input. Data types supported: U8/QASYMM8/QASYMM8_SIGNED/S16/F16/F32 - * @param[in] output Output tensor. Data types supported: U8/QASYMM8/QASYMM8_SIGNED/S16/F16/F32 - * @param[in] policy Policy to use to handle overflow. Convert policy cannot be WRAP if datatype is quantized. + * @param[in] input1 First tensor input. Data types supported: U8/QASYMM8/QASYMM8_SIGNED/S16/F16/F32 + * @param[in] input2 Second tensor input. Data types supported: U8/QASYMM8/QASYMM8_SIGNED/S16/F16/F32 + * @param[in] output Output tensor. Data types supported: U8/QASYMM8/QASYMM8_SIGNED/S16/F16/F32 + * @param[in] policy Policy to use to handle overflow. Convert policy cannot be WRAP if datatype is quantized. + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. Currently not supported. * * @return a status */ - static Status validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, ConvertPolicy policy); + static Status validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, ConvertPolicy policy, const ActivationLayerInfo &act_info = ActivationLayerInfo()); }; } #endif /* ARM_COMPUTE_NEARITHMETICSUBTRACTION_H */ diff --git a/arm_compute/runtime/NEON/functions/NEElementwiseOperations.h b/arm_compute/runtime/NEON/functions/NEElementwiseOperations.h index e5af6bc841..cac105cdb9 100644 --- a/arm_compute/runtime/NEON/functions/NEElementwiseOperations.h +++ b/arm_compute/runtime/NEON/functions/NEElementwiseOperations.h @@ -41,20 +41,22 @@ class NEElementwiseMax : public INESimpleFunction public: /** Initialise the kernel's inputs, output and conversion policy. * - * @param[in, out] input1 First tensor input. Data types supported: QASYMM8/QASYMM8_SIGNED/S16/F16/S32/F32. - * @param[in, out] input2 Second tensor input. Data types supported: Same as @p input1. - * @param[out] output Output tensor. Data types supported: Same as @p input1. + * @param[in, out] input1 First tensor input. Data types supported: QASYMM8/QASYMM8_SIGNED/S16/F16/S32/F32. + * @param[in, out] input2 Second tensor input. Data types supported: Same as @p input1. + * @param[out] output Output tensor. Data types supported: Same as @p input1. + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. Currently not supported. */ - void configure(ITensor *input1, ITensor *input2, ITensor *output); + void configure(ITensor *input1, ITensor *input2, ITensor *output, const ActivationLayerInfo &act_info = ActivationLayerInfo()); /** Static function to check if given info will lead to a valid configuration of @ref NEArithmeticOperationKernel for max * - * @param[in] input1 First tensor input info. Data types supported: QASYMM8/QASYMM8_SIGNED/S16/F16/S32/F32. - * @param[in] input2 Second tensor input info. Data types supported: Same as @p input1. - * @param[in] output Output tensor info. Data types supported: Same as @p input1. + * @param[in] input1 First tensor input info. Data types supported: QASYMM8/QASYMM8_SIGNED/S16/F16/S32/F32. + * @param[in] input2 Second tensor input info. Data types supported: Same as @p input1. + * @param[in] output Output tensor info. Data types supported: Same as @p input1. + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. Currently not supported. * * @return a status */ - static Status validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output); + static Status validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, const ActivationLayerInfo &act_info = ActivationLayerInfo()); }; /** Basic function to run @ref NEArithmeticOperationKernel for min @@ -67,20 +69,22 @@ class NEElementwiseMin : public INESimpleFunction public: /** Initialise the kernel's inputs, output and conversion policy. * - * @param[in, out] input1 First tensor input. Data types supported: QASYMM8/QASYMM8_SIGNED/S16/F16/S32/F32. - * @param[in, out] input2 Second tensor input. Data types supported: Same as @p input1. - * @param[out] output Output tensor. Data types supported: Same as @p input1. + * @param[in, out] input1 First tensor input. Data types supported: QASYMM8/QASYMM8_SIGNED/S16/F16/S32/F32. + * @param[in, out] input2 Second tensor input. Data types supported: Same as @p input1. + * @param[out] output Output tensor. Data types supported: Same as @p input1. + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. Currently not supported. */ - void configure(ITensor *input1, ITensor *input2, ITensor *output); + void configure(ITensor *input1, ITensor *input2, ITensor *output, const ActivationLayerInfo &act_info = ActivationLayerInfo()); /** Static function to check if given info will lead to a valid configuration of @ref NEArithmeticOperationKernel for min * - * @param[in] input1 First tensor input info. Data types supported: QASYMM8/QASYMM8_SIGNED/S16/F16/S32/F32. - * @param[in] input2 Second tensor input info. Data types supported: Same as @p input1. - * @param[in] output Output tensor info. Data types supported: Same as @p input1. + * @param[in] input1 First tensor input info. Data types supported: QASYMM8/QASYMM8_SIGNED/S16/F16/S32/F32. + * @param[in] input2 Second tensor input info. Data types supported: Same as @p input1. + * @param[in] output Output tensor info. Data types supported: Same as @p input1. + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. Currently not supported. * * @return a status */ - static Status validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output); + static Status validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, const ActivationLayerInfo &act_info = ActivationLayerInfo()); }; /** Basic function to run @ref NEArithmeticOperationKernel for squared difference @@ -93,20 +97,22 @@ class NEElementwiseSquaredDiff : public INESimpleFunction public: /** Initialise the kernel's inputs, output and conversion policy. * - * @param[in, out] input1 First tensor input. Data types supported: QASYMM8/QASYMM8_SIGNED/S16/F16/S32/F32. - * @param[in, out] input2 Second tensor input. Data types supported: Same as @p input1. - * @param[out] output Output tensor. Data types supported: Same as @p input1. + * @param[in, out] input1 First tensor input. Data types supported: QASYMM8/QASYMM8_SIGNED/S16/F16/S32/F32. + * @param[in, out] input2 Second tensor input. Data types supported: Same as @p input1. + * @param[out] output Output tensor. Data types supported: Same as @p input1. + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. Currently not supported. */ - void configure(ITensor *input1, ITensor *input2, ITensor *output); + void configure(ITensor *input1, ITensor *input2, ITensor *output, const ActivationLayerInfo &act_info = ActivationLayerInfo()); /** Static function to check if given info will lead to a valid configuration of @ref NEArithmeticOperationKernel for squared difference * - * @param[in] input1 First tensor input info. Data types supported: QASYMM8/QASYMM8_SIGNED/S16/F16/S32/F32. - * @param[in] input2 Second tensor input info. Data types supported: Same as @p input1. - * @param[in] output Output tensor info. Data types supported: Same as @p input1. + * @param[in] input1 First tensor input info. Data types supported: QASYMM8/QASYMM8_SIGNED/S16/F16/S32/F32. + * @param[in] input2 Second tensor input info. Data types supported: Same as @p input1. + * @param[in] output Output tensor info. Data types supported: Same as @p input1. + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. Currently not supported. * * @return a status */ - static Status validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output); + static Status validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, const ActivationLayerInfo &act_info = ActivationLayerInfo()); }; /** Basic function to run @ref NEArithmeticOperationKernel for division @@ -119,20 +125,22 @@ class NEElementwiseDivision : public INESimpleFunction public: /** Initialise the kernel's inputs, output and conversion policy. * - * @param[in, out] input1 First tensor input. Data types supported: F16/F32. - * @param[in, out] input2 Second tensor input. Data types supported: Same as @p input1. - * @param[out] output Output tensor. Data types supported: Same as @p input1. + * @param[in, out] input1 First tensor input. Data types supported: F16/F32. + * @param[in, out] input2 Second tensor input. Data types supported: Same as @p input1. + * @param[out] output Output tensor. Data types supported: Same as @p input1. + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. Currently not supported. */ - void configure(ITensor *input1, ITensor *input2, ITensor *output); + void configure(ITensor *input1, ITensor *input2, ITensor *output, const ActivationLayerInfo &act_info = ActivationLayerInfo()); /** Static function to check if given info will lead to a valid configuration of @ref NEArithmeticOperationKernel for division * - * @param[in] input1 First tensor input info. Data types supported: F16/F32. - * @param[in] input2 Second tensor input info. Data types supported: Same as @p input1. - * @param[in] output Output tensor info. Data types supported: Same as @p input1. + * @param[in] input1 First tensor input info. Data types supported: F16/F32. + * @param[in] input2 Second tensor input info. Data types supported: Same as @p input1. + * @param[in] output Output tensor info. Data types supported: Same as @p input1. + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. Currently not supported. * * @return a status */ - static Status validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output); + static Status validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, const ActivationLayerInfo &act_info = ActivationLayerInfo()); }; /** Basic function to run @ref NEArithmeticOperationKernel for power @@ -146,20 +154,22 @@ class NEElementwisePower : public INESimpleFunction public: /** Initialise the kernel's inputs, output and conversion policy. * - * @param[in, out] input1 First tensor input. Data types supported: F16/F32. - * @param[in, out] input2 Second tensor input. Data types supported: Same as @p input1. - * @param[out] output Output tensor. Data types supported: Same as @p input1. + * @param[in, out] input1 First tensor input. Data types supported: F16/F32. + * @param[in, out] input2 Second tensor input. Data types supported: Same as @p input1. + * @param[out] output Output tensor. Data types supported: Same as @p input1. + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. Currently not supported. */ - void configure(ITensor *input1, ITensor *input2, ITensor *output); + void configure(ITensor *input1, ITensor *input2, ITensor *output, const ActivationLayerInfo &act_info = ActivationLayerInfo()); /** Static function to check if given info will lead to a valid configuration of @ref NEArithmeticOperationKernel for power * - * @param[in] input1 First tensor input info. Data types supported: F16/F32. - * @param[in] input2 Second tensor input info. Data types supported: Same as @p input1. - * @param[in] output Output tensor info. Data types supported: Same as @p input1. + * @param[in] input1 First tensor input info. Data types supported: F16/F32. + * @param[in] input2 Second tensor input info. Data types supported: Same as @p input1. + * @param[in] output Output tensor info. Data types supported: Same as @p input1. + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. Currently not supported. * * @return a status */ - static Status validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output); + static Status validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, const ActivationLayerInfo &act_info = ActivationLayerInfo()); }; /** Basic function to run @ref NEComparisonOperationKernel. diff --git a/arm_compute/runtime/NEON/functions/NEPixelWiseMultiplication.h b/arm_compute/runtime/NEON/functions/NEPixelWiseMultiplication.h index ede4327bfb..2b31032931 100644 --- a/arm_compute/runtime/NEON/functions/NEPixelWiseMultiplication.h +++ b/arm_compute/runtime/NEON/functions/NEPixelWiseMultiplication.h @@ -57,8 +57,10 @@ public: * Scale must be positive and its value must be either 1/255 or 1/2^n where n is between 0 and 15. * @param[in] overflow_policy Overflow policy. ConvertPolicy cannot be WRAP if datatype is QASYMM8, QASYMM8_SIGNED or QSYMM16. * @param[in] rounding_policy Rounding policy. + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. Currently not supported. */ - void configure(ITensor *input1, ITensor *input2, ITensor *output, float scale, ConvertPolicy overflow_policy, RoundingPolicy rounding_policy); + void configure(ITensor *input1, ITensor *input2, ITensor *output, float scale, ConvertPolicy overflow_policy, RoundingPolicy rounding_policy, + const ActivationLayerInfo &act_info = ActivationLayerInfo()); /** Static function to check if given info will lead to a valid configuration of @ref NEPixelWiseMultiplication * * @note For @p scale equal to 1/255 only round to nearest even (implemented as round half up) is supported. @@ -79,10 +81,12 @@ public: * Scale must be positive and its value must be either 1/255 or 1/2^n where n is between 0 and 15. * @param[in] overflow_policy Overflow policy. ConvertPolicy cannot be WRAP if datatype is QASYMM8, QASYMM8_SIGNED or QSYMM16. * @param[in] rounding_policy Rounding policy. + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. Currently not supported. * * @return a status */ - static Status validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, float scale, ConvertPolicy overflow_policy, RoundingPolicy rounding_policy); + static Status validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, float scale, ConvertPolicy overflow_policy, RoundingPolicy rounding_policy, + const ActivationLayerInfo &act_info = ActivationLayerInfo()); }; /** Basic function to run @ref NEComplexPixelWiseMultiplicationKernel. */ @@ -91,20 +95,22 @@ class NEComplexPixelWiseMultiplication : public INESimpleFunction public: /** Initialise the kernel's inputs, output. * - * @param[in, out] input1 An input tensor. Data types supported: F32. Number of channels supported: 2 (complex tensor). - * The input tensor is [in, out] because its TensorInfo might be modified inside the kernel in case of broadcasting of dimension 0. - * @param[in, out] input2 An input tensor. Data types supported: same as @p input1. Number of channels supported: same as @p input1. - * The input tensor is [in, out] because its TensorInfo might be modified inside the kernel in case of broadcasting of dimension 0. - * @param[out] output The output tensor. Data types supported: same as @p input1. Number of channels: same as @p input1. + * @param[in, out] input1 An input tensor. Data types supported: F32. Number of channels supported: 2 (complex tensor). + * The input tensor is [in, out] because its TensorInfo might be modified inside the kernel in case of broadcasting of dimension 0. + * @param[in, out] input2 An input tensor. Data types supported: same as @p input1. Number of channels supported: same as @p input1. + * The input tensor is [in, out] because its TensorInfo might be modified inside the kernel in case of broadcasting of dimension 0. + * @param[out] output The output tensor. Data types supported: same as @p input1. Number of channels: same as @p input1. + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. Currently not supported. */ - void configure(ITensor *input1, ITensor *input2, ITensor *output); + void configure(ITensor *input1, ITensor *input2, ITensor *output, const ActivationLayerInfo &act_info = ActivationLayerInfo()); /** Static function to check if given info will lead to a valid configuration of @ref NEComplexPixelWiseMultiplication * - * @param[in] input1 An input tensor info. Data types supported: F32. Number of channels supported: 2 (complex tensor). - * @param[in] input2 An input tensor info. Data types supported: same as @p input1. Number of channels supported: same as @p input1. - * @param[in] output The output tensor info. Data types supported: same as @p input1. Number of channels supported: same as @p input1. + * @param[in] input1 An input tensor info. Data types supported: F32. Number of channels supported: 2 (complex tensor). + * @param[in] input2 An input tensor info. Data types supported: same as @p input1. Number of channels supported: same as @p input1. + * @param[in] output The output tensor info. Data types supported: same as @p input1. Number of channels supported: same as @p input1. + * @param[in] act_info (Optional) Activation layer information in case of a fused activation. Currently not supported. */ - static Status validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output); + static Status validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, const ActivationLayerInfo &act_info = ActivationLayerInfo()); }; } #endif /*ARM_COMPUTE_NEPIXELWISEMULTIPLICATION_H */ diff --git a/src/core/CL/cl_kernels/activation_float_helpers.h b/src/core/CL/cl_kernels/activation_float_helpers.h index 8590f25635..a1e742da0d 100644 --- a/src/core/CL/cl_kernels/activation_float_helpers.h +++ b/src/core/CL/cl_kernels/activation_float_helpers.h @@ -72,6 +72,6 @@ // Identity Activation #define identity_op(DATA_TYPE, x, A_VAL, B_VAL) (x) -#define OP(op, DATA_TYPE, x, A_VAL, B_VAL) op##_op(DATA_TYPE, x, A_VAL, B_VAL) +#define ACT_OP(op, DATA_TYPE, x, A_VAL, B_VAL) op##_op(DATA_TYPE, x, A_VAL, B_VAL) -#define ACTIVATION(op, DATA_TYPE, x, A_VAL, B_VAL) OP(op, DATA_TYPE, x, A_VAL, B_VAL) +#define ACTIVATION(op, DATA_TYPE, x, A_VAL, B_VAL) ACT_OP(op, DATA_TYPE, x, A_VAL, B_VAL) diff --git a/src/core/CL/cl_kernels/elementwise_operation.cl b/src/core/CL/cl_kernels/elementwise_operation.cl index 42d6d33e03..9b87b526f7 100644 --- a/src/core/CL/cl_kernels/elementwise_operation.cl +++ b/src/core/CL/cl_kernels/elementwise_operation.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2019 ARM Limited. + * Copyright (c) 2018-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -44,6 +44,11 @@ #define OP_FUN_NAME(op) OP_FUN_NAME_STR(op) #if defined(OP) && defined(DATA_TYPE_IN1) && defined(DATA_TYPE_IN2) && defined(DATA_TYPE_OUT) && defined(VEC_SIZE) + +#if defined(ACTIVATION_TYPE) +#include "activation_float_helpers.h" +#endif // defined(ACTIVATION_TYPE) + /** This function executes an element-wise operation among two tensors. * * @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,7 +99,12 @@ __kernel void OP_FUN_NAME(OP)( in_b = CONVERT(VLOAD(VEC_SIZE)(0, (__global DATA_TYPE_IN2 *)in2.ptr), VEC_DATA_TYPE(DATA_TYPE_OUT, VEC_SIZE)); // Calculate and store result +#if defined(ACTIVATION_TYPE) + VSTORE(VEC_SIZE) + (ACTIVATION(ACTIVATION_TYPE, DATA_TYPE_OUT, CONVERT(OP(in_a, in_b), VEC_DATA_TYPE(DATA_TYPE_OUT, VEC_SIZE)), A_VAL, B_VAL), 0, (__global DATA_TYPE_OUT *)out.ptr); +#else // defined(ACTIVATION_TYPE) VSTORE(VEC_SIZE) (OP(in_a, in_b), 0, (__global DATA_TYPE_OUT *)out.ptr); +#endif // defined(ACTIVATION_TYPE) } #endif /* defined(DATA_TYPE_IN1) && defined(DATA_TYPE_IN2) && defined(DATA_TYPE_OUT) && defined(VEC_SIZE) */ diff --git a/src/core/CL/cl_kernels/pixelwise_mul_float.cl b/src/core/CL/cl_kernels/pixelwise_mul_float.cl index d0e04b2ffe..aad4becc1a 100644 --- a/src/core/CL/cl_kernels/pixelwise_mul_float.cl +++ b/src/core/CL/cl_kernels/pixelwise_mul_float.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016-2019 ARM Limited. + * Copyright (c) 2016-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -31,6 +31,11 @@ #define CONVERT_OP_FLOAT(x, type, round) CONVERT_OP_FLOAT_STR(x, type, round) #if defined(DATA_TYPE_IN1) && defined(DATA_TYPE_IN2) && defined(DATA_TYPE_RES) && defined(DATA_TYPE_OUT) + +#if defined(ACTIVATION_TYPE) +#include "activation_float_helpers.h" +#endif // defined(ACTIVATION_TYPE) + /** Performs a pixelwise multiplication with float scale of either integer or float inputs. * * @attention The inputs and output data types need to be passed at compile time using -DDATA_TYPE_IN1, -DDATA_TYPE_IN2 and -DDATA_TYPE_OUT: @@ -91,8 +96,12 @@ __kernel void pixelwise_mul_float( res = CONVERT_OP_FLOAT(CONVERT_OP_FLOAT((convert_float16(in1_data * in2_data) * scale), VEC_DATA_TYPE(DATA_TYPE_RES, 16), ROUND), VEC_DATA_TYPE(DATA_TYPE_OUT, 16), ROUND); #endif /* DATA_TYPE_FLOAT */ +#if defined(ACTIVATION_TYPE) + vstore16(ACTIVATION(ACTIVATION_TYPE, DATA_TYPE_OUT, res, A_VAL, B_VAL), 0, (__global DATA_TYPE_OUT *)out.ptr); +#else // defined(ACTIVATION_TYPE) // Store result vstore16(res, 0, (__global DATA_TYPE_OUT *)out.ptr); +#endif // defined(ACTIVATION_TYPE) } #endif /* defined(DATA_TYPE_IN1) && defined(DATA_TYPE_IN2) && defined(DATA_TYPE_RES) && defined(DATA_TYPE_OUT) */ @@ -140,6 +149,10 @@ __kernel void pixelwise_mul_complex( // Perform complex multiplication float2 res = { vin1.x *vin2.x - vin1.y * vin2.y, vin1.x *vin2.y + vin2.x * vin1.y }; +#if defined(ACTIVATION_TYPE) + vstore2(ACTIVATION(ACTIVATION_TYPE, float, res, A_VAL, B_VAL), 0, (__global float *)out.ptr); +#else // defined(ACTIVATION_TYPE) // Store result vstore2(res, 0, (__global float *)out.ptr); +#endif // defined(ACTIVATION_TYPE) } diff --git a/src/core/CL/kernels/CLElementwiseOperationKernel.cpp b/src/core/CL/kernels/CLElementwiseOperationKernel.cpp index 1ac35a286f..0f2e26f186 100644 --- a/src/core/CL/kernels/CLElementwiseOperationKernel.cpp +++ b/src/core/CL/kernels/CLElementwiseOperationKernel.cpp @@ -231,7 +231,7 @@ std::pair validate_and_configure_window_for_division(ITensorInfo } // namespace CLElementwiseOperationKernel::CLElementwiseOperationKernel() - : _input1(nullptr), _input2(nullptr), _output(nullptr) + : _act_info(), _input1(nullptr), _input2(nullptr), _output(nullptr) { } @@ -256,6 +256,12 @@ void CLElementwiseOperationKernel::configure_common(const ICLTensor *input1, con // Set kernel build options CLBuildOptions build_opts = generate_build_options(*input1->info(), *input2->info(), *output->info()); + if(_act_info.enabled()) + { + build_opts.add_option("-DACTIVATION_TYPE=" + lower_string(string_from_activation_func(_act_info.activation()))); + build_opts.add_option("-DA_VAL=" + float_to_string_with_full_precision(_act_info.a())); + build_opts.add_option("-DB_VAL=" + float_to_string_with_full_precision(_act_info.b())); + } // Create kernel _kernel = static_cast(CLKernelLibrary::get().create_kernel(kernel_name, build_opts.options())); @@ -320,19 +326,23 @@ BorderSize CLElementwiseOperationKernel::border_size() const /** Arithmetic operations with saturation*/ -void CLSaturatedArithmeticOperationKernel::configure(ArithmeticOperation op, const ICLTensor *input1, const ICLTensor *input2, ICLTensor *output, const ConvertPolicy &policy) +void CLSaturatedArithmeticOperationKernel::configure(ArithmeticOperation op, const ICLTensor *input1, const ICLTensor *input2, ICLTensor *output, const ConvertPolicy &policy, + const ActivationLayerInfo &act_info) { - _policy = policy; - _op = op; + _policy = policy; + _op = op; + _act_info = act_info; configure_common(input1, input2, output); } -Status CLSaturatedArithmeticOperationKernel::validate(ArithmeticOperation op, const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, const ConvertPolicy &policy) +Status CLSaturatedArithmeticOperationKernel::validate(ArithmeticOperation op, const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, const ConvertPolicy &policy, + const ActivationLayerInfo &act_info) { ARM_COMPUTE_UNUSED(op, policy); ARM_COMPUTE_RETURN_ERROR_ON_NULLPTR(input1, input2, output); ARM_COMPUTE_RETURN_ON_ERROR(validate_arguments_with_arithmetic_rules(*input1, *input2, *output)); ARM_COMPUTE_RETURN_ON_ERROR(validate_and_configure_window_for_arithmetic_operators(*input1->clone(), *input2->clone(), *output->clone()).first); + ARM_COMPUTE_RETURN_ERROR_ON(act_info.enabled() && !is_data_type_float(output->data_type())); return Status{}; } @@ -369,13 +379,14 @@ std::string CLSaturatedArithmeticOperationKernel::name() /** Arithmetic operations*/ -void CLArithmeticOperationKernel::configure(ArithmeticOperation op, const ICLTensor *input1, const ICLTensor *input2, ICLTensor *output) +void CLArithmeticOperationKernel::configure(ArithmeticOperation op, const ICLTensor *input1, const ICLTensor *input2, ICLTensor *output, const ActivationLayerInfo &act_info) { - _op = op; + _op = op; + _act_info = act_info; configure_common(input1, input2, output); } -Status CLArithmeticOperationKernel::validate(ArithmeticOperation op, const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output) +Status CLArithmeticOperationKernel::validate(ArithmeticOperation op, const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, const ActivationLayerInfo &act_info) { ARM_COMPUTE_RETURN_ERROR_ON_NULLPTR(input1, input2, output); if(op == ArithmeticOperation::DIV || op == ArithmeticOperation::POWER) @@ -389,6 +400,7 @@ Status CLArithmeticOperationKernel::validate(ArithmeticOperation op, const ITens ARM_COMPUTE_RETURN_ON_ERROR(validate_arguments_with_arithmetic_rules(*input1, *input2, *output)); ARM_COMPUTE_RETURN_ON_ERROR(validate_and_configure_window_for_arithmetic_operators(*input1->clone(), *input2->clone(), *output->clone()).first); } + ARM_COMPUTE_RETURN_ERROR_ON(act_info.enabled() && !is_data_type_float(output->data_type())); return Status{}; } diff --git a/src/core/CL/kernels/CLPixelWiseMultiplicationKernel.cpp b/src/core/CL/kernels/CLPixelWiseMultiplicationKernel.cpp index d31c1de402..ff5afa3d95 100644 --- a/src/core/CL/kernels/CLPixelWiseMultiplicationKernel.cpp +++ b/src/core/CL/kernels/CLPixelWiseMultiplicationKernel.cpp @@ -46,7 +46,7 @@ namespace constexpr unsigned int num_elems_processed_per_iteration = 16; Status validate_arguments(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, float scale, - ConvertPolicy overflow_policy, RoundingPolicy rounding_policy) + ConvertPolicy overflow_policy, RoundingPolicy rounding_policy, const ActivationLayerInfo &act_info) { ARM_COMPUTE_UNUSED(overflow_policy); ARM_COMPUTE_UNUSED(rounding_policy); @@ -64,6 +64,7 @@ Status validate_arguments(const ITensorInfo *input1, const ITensorInfo *input2, DataType::S16, DataType::QSYMM16, DataType::F16, DataType::F32); ARM_COMPUTE_RETURN_ERROR_ON_MSG(scale < 0, "Scale cannot be negative."); + ARM_COMPUTE_RETURN_ERROR_ON(act_info.enabled() && !is_data_type_float(output->data_type())); const TensorShape &out_shape = TensorShape::broadcast_shape(input1->tensor_shape(), input2->tensor_shape()); @@ -148,11 +149,11 @@ CLPixelWiseMultiplicationKernel::CLPixelWiseMultiplicationKernel() } void CLPixelWiseMultiplicationKernel::configure(const ICLTensor *input1, const ICLTensor *input2, ICLTensor *output, float scale, - ConvertPolicy overflow_policy, RoundingPolicy rounding_policy) + ConvertPolicy overflow_policy, RoundingPolicy rounding_policy, const ActivationLayerInfo &act_info) { ARM_COMPUTE_ERROR_ON_NULLPTR(input1, input2, output); ARM_COMPUTE_ERROR_THROW_ON(validate_arguments(input1->info(), input2->info(), output->info(), - scale, overflow_policy, rounding_policy)); + scale, overflow_policy, rounding_policy, act_info)); // Configure kernel window auto win_config = validate_and_configure_window(input1->info(), input2->info(), output->info()); @@ -227,6 +228,12 @@ void CLPixelWiseMultiplicationKernel::configure(const ICLTensor *input1, const I build_opts.add_option_if_else(overflow_policy == ConvertPolicy::WRAP || is_data_type_float(output->info()->data_type()), "-DWRAP", "-DSATURATE"); build_opts.add_option_if_else(rounding_policy == RoundingPolicy::TO_ZERO, "-DROUND=_rtz", "-DROUND=_rte"); build_opts.add_option("-DDATA_TYPE_RES=" + compute_type); + if(act_info.enabled()) + { + build_opts.add_option("-DACTIVATION_TYPE=" + lower_string(string_from_activation_func(act_info.activation()))); + build_opts.add_option("-DA_VAL=" + float_to_string_with_full_precision(act_info.a())); + build_opts.add_option("-DB_VAL=" + float_to_string_with_full_precision(act_info.b())); + } } // Create kernel @@ -248,10 +255,10 @@ void CLPixelWiseMultiplicationKernel::configure(const ICLTensor *input1, const I } Status CLPixelWiseMultiplicationKernel::validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, float scale, - ConvertPolicy overflow_policy, RoundingPolicy rounding_policy) + ConvertPolicy overflow_policy, RoundingPolicy rounding_policy, const ActivationLayerInfo &act_info) { ARM_COMPUTE_ERROR_ON_NULLPTR(input1, input2, output); - ARM_COMPUTE_RETURN_ON_ERROR(validate_arguments(input1, input2, output, scale, overflow_policy, rounding_policy)); + ARM_COMPUTE_RETURN_ON_ERROR(validate_arguments(input1, input2, output, scale, overflow_policy, rounding_policy, act_info)); ARM_COMPUTE_RETURN_ON_ERROR(validate_and_configure_window(input1->clone().get(), input2->clone().get(), output->clone().get()).first); return Status{}; @@ -311,7 +318,7 @@ namespace { constexpr unsigned int num_elems_processed_per_iteration_complex = 1; -Status validate_arguments_complex(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output) +Status validate_arguments_complex(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, const ActivationLayerInfo &act_info) { ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input1, 2, DataType::F32); ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input2, 2, DataType::F32); @@ -319,6 +326,7 @@ Status validate_arguments_complex(const ITensorInfo *input1, const ITensorInfo * const TensorShape &out_shape = TensorShape::broadcast_shape(input1->tensor_shape(), input2->tensor_shape()); ARM_COMPUTE_RETURN_ERROR_ON_MSG(out_shape.total_size() == 0, "Inputs are not broadcast compatible"); + ARM_COMPUTE_RETURN_ERROR_ON(act_info.enabled() && !is_data_type_float(output->data_type())); // Validate in case of configured output if(output->total_size() > 0) @@ -364,10 +372,10 @@ CLComplexPixelWiseMultiplicationKernel::CLComplexPixelWiseMultiplicationKernel() { } -void CLComplexPixelWiseMultiplicationKernel::configure(const ICLTensor *input1, const ICLTensor *input2, ICLTensor *output) +void CLComplexPixelWiseMultiplicationKernel::configure(const ICLTensor *input1, const ICLTensor *input2, ICLTensor *output, const ActivationLayerInfo &act_info) { ARM_COMPUTE_ERROR_ON_NULLPTR(input1, input2, output); - ARM_COMPUTE_ERROR_THROW_ON(validate_arguments_complex(input1->info(), input2->info(), output->info())); + ARM_COMPUTE_ERROR_THROW_ON(validate_arguments_complex(input1->info(), input2->info(), output->info(), act_info)); // Configure kernel window auto win_config = validate_and_configure_window_complex(input1->info(), input2->info(), output->info()); @@ -377,16 +385,24 @@ void CLComplexPixelWiseMultiplicationKernel::configure(const ICLTensor *input1, _input2 = input2; _output = output; + CLBuildOptions build_opts; + if(act_info.enabled()) + { + build_opts.add_option("-DACTIVATION_TYPE=" + lower_string(string_from_activation_func(act_info.activation()))); + build_opts.add_option("-DA_VAL=" + float_to_string_with_full_precision(act_info.a())); + build_opts.add_option("-DB_VAL=" + float_to_string_with_full_precision(act_info.b())); + } + // Create kernel - _kernel = static_cast(CLKernelLibrary::get().create_kernel("pixelwise_mul_complex")); + _kernel = static_cast(CLKernelLibrary::get().create_kernel("pixelwise_mul_complex", build_opts.options())); ICLKernel::configure_internal(win_config.second); } -Status CLComplexPixelWiseMultiplicationKernel::validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output) +Status CLComplexPixelWiseMultiplicationKernel::validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, const ActivationLayerInfo &act_info) { ARM_COMPUTE_ERROR_ON_NULLPTR(input1, input2, output); - ARM_COMPUTE_RETURN_ON_ERROR(validate_arguments_complex(input1, input2, output)); + ARM_COMPUTE_RETURN_ON_ERROR(validate_arguments_complex(input1, input2, output, act_info)); ARM_COMPUTE_RETURN_ON_ERROR(validate_and_configure_window_complex(input1->clone().get(), input2->clone().get(), output->clone().get()).first); return Status{}; diff --git a/src/graph/mutators/NodeFusionMutator.cpp b/src/graph/mutators/NodeFusionMutator.cpp index 273e6ce746..ae53b8ff75 100644 --- a/src/graph/mutators/NodeFusionMutator.cpp +++ b/src/graph/mutators/NodeFusionMutator.cpp @@ -294,13 +294,20 @@ IGraphMutator::MutationType NodeFusionMutator::type() const void NodeFusionMutator::mutate(Graph &g) { // Supported activations when fusing - const std::set supported_fused_activations = { Activation::RELU, Activation::BOUNDED_RELU, Activation::LU_BOUNDED_RELU }; + const std::set supported_fused_activations_conv = { Activation::RELU, Activation::BOUNDED_RELU, Activation::LU_BOUNDED_RELU }; + const std::set supported_fused_activations_eltwise = { Activation::RELU, Activation::BOUNDED_RELU, Activation::LU_BOUNDED_RELU, + Activation::TANH, Activation::LOGISTIC + }; // Preconditions auto empty_prec = [](INode &) { return true; }; + auto cl_target_prec = [](INode & n) + { + return n.assigned_target() == Target::CL; + }; auto qs8_prec = [&g](INode & n) { ARM_COMPUTE_ERROR_ON(n.output(0) == nullptr); @@ -315,10 +322,11 @@ void NodeFusionMutator::mutate(Graph &g) }; // Fusion mutations - detail::fuse_layer(g, empty_prec, detail::fuse_node_with_activation, supported_fused_activations); - detail::fuse_layer(g, empty_prec, detail::fuse_node_with_activation, supported_fused_activations); - detail::fuse_layer(g, qs8_prec, detail::fuse_node_with_activation, supported_fused_activations); - detail::fuse_layer(g, empty_prec, detail::fuse_node_with_activation, supported_fused_activations); + detail::fuse_layer(g, empty_prec, detail::fuse_node_with_activation, supported_fused_activations_conv); + detail::fuse_layer(g, empty_prec, detail::fuse_node_with_activation, supported_fused_activations_conv); + detail::fuse_layer(g, qs8_prec, detail::fuse_node_with_activation, supported_fused_activations_conv); + detail::fuse_layer(g, empty_prec, detail::fuse_node_with_activation, supported_fused_activations_conv); + detail::fuse_layer(g, cl_target_prec, detail::fuse_node_with_activation, supported_fused_activations_eltwise); detail::fuse_layer(g, empty_prec, detail::fuse_convolution_with_batch_normalization); detail::fuse_layer(g, empty_prec, detail::fuse_depthwise_convolution_with_batch_normalization); } diff --git a/src/graph/nodes/EltwiseLayerNode.cpp b/src/graph/nodes/EltwiseLayerNode.cpp index a83a5fb3b2..92d183e693 100644 --- a/src/graph/nodes/EltwiseLayerNode.cpp +++ b/src/graph/nodes/EltwiseLayerNode.cpp @@ -52,6 +52,16 @@ RoundingPolicy EltwiseLayerNode::rounding_policy() const return descriptor.r_policy; } +ActivationLayerInfo EltwiseLayerNode::fused_activation() const +{ + return descriptor.fused_activation; +} + +void EltwiseLayerNode::set_fused_activation(ActivationLayerInfo fused_activation) +{ + descriptor.fused_activation = fused_activation; +} + bool EltwiseLayerNode::forward_descriptors() { if((input_id(0) != NullTensorID) && (output_id(0) != NullTensorID)) diff --git a/src/runtime/CL/functions/CLElementwiseOperations.cpp b/src/runtime/CL/functions/CLElementwiseOperations.cpp index 69cebc7180..7636a87e93 100644 --- a/src/runtime/CL/functions/CLElementwiseOperations.cpp +++ b/src/runtime/CL/functions/CLElementwiseOperations.cpp @@ -47,96 +47,96 @@ void configure_border_handler(CLFillBorderKernel &border_handler, BorderSize bor } } // namespace -void CLArithmeticAddition::configure(ICLTensor *input1, ICLTensor *input2, ICLTensor *output, ConvertPolicy policy) +void CLArithmeticAddition::configure(ICLTensor *input1, ICLTensor *input2, ICLTensor *output, ConvertPolicy policy, const ActivationLayerInfo &act_info) { auto k = arm_compute::support::cpp14::make_unique(); - k->configure(ArithmeticOperation::ADD, input1, input2, output, policy); + k->configure(ArithmeticOperation::ADD, input1, input2, output, policy, act_info); _kernel = std::move(k); configure_border_handler(_border_handler, _kernel->border_size(), input1, input2, output); } -Status CLArithmeticAddition::validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, ConvertPolicy policy) +Status CLArithmeticAddition::validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, ConvertPolicy policy, const ActivationLayerInfo &act_info) { - return CLSaturatedArithmeticOperationKernel::validate(ArithmeticOperation::ADD, input1, input2, output, policy); + return CLSaturatedArithmeticOperationKernel::validate(ArithmeticOperation::ADD, input1, input2, output, policy, act_info); } -void CLArithmeticSubtraction::configure(ICLTensor *input1, ICLTensor *input2, ICLTensor *output, ConvertPolicy policy) +void CLArithmeticSubtraction::configure(ICLTensor *input1, ICLTensor *input2, ICLTensor *output, ConvertPolicy policy, const ActivationLayerInfo &act_info) { auto k = arm_compute::support::cpp14::make_unique(); - k->configure(ArithmeticOperation::SUB, input1, input2, output, policy); + k->configure(ArithmeticOperation::SUB, input1, input2, output, policy, act_info); _kernel = std::move(k); configure_border_handler(_border_handler, _kernel->border_size(), input1, input2, output); } -Status CLArithmeticSubtraction::validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, ConvertPolicy policy) +Status CLArithmeticSubtraction::validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, ConvertPolicy policy, const ActivationLayerInfo &act_info) { ARM_COMPUTE_UNUSED(policy); - return CLSaturatedArithmeticOperationKernel::validate(ArithmeticOperation::SUB, input1, input2, output, policy); + return CLSaturatedArithmeticOperationKernel::validate(ArithmeticOperation::SUB, input1, input2, output, policy, act_info); } -void CLArithmeticDivision::configure(ICLTensor *input1, ICLTensor *input2, ICLTensor *output) +void CLArithmeticDivision::configure(ICLTensor *input1, ICLTensor *input2, ICLTensor *output, const ActivationLayerInfo &act_info) { auto k = arm_compute::support::cpp14::make_unique(); - k->configure(ArithmeticOperation::DIV, input1, input2, output); + k->configure(ArithmeticOperation::DIV, input1, input2, output, act_info); _kernel = std::move(k); configure_border_handler(_border_handler, _kernel->border_size(), input1, input2, output); } -Status CLArithmeticDivision::validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output) +Status CLArithmeticDivision::validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, const ActivationLayerInfo &act_info) { - return CLArithmeticOperationKernel::validate(ArithmeticOperation::DIV, input1, input2, output); + return CLArithmeticOperationKernel::validate(ArithmeticOperation::DIV, input1, input2, output, act_info); } -void CLElementwiseMax::configure(ICLTensor *input1, ICLTensor *input2, ICLTensor *output) +void CLElementwiseMax::configure(ICLTensor *input1, ICLTensor *input2, ICLTensor *output, const ActivationLayerInfo &act_info) { auto k = arm_compute::support::cpp14::make_unique(); - k->configure(ArithmeticOperation::MAX, input1, input2, output); + k->configure(ArithmeticOperation::MAX, input1, input2, output, act_info); _kernel = std::move(k); configure_border_handler(_border_handler, _kernel->border_size(), input1, input2, output); } -Status CLElementwiseMax::validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output) +Status CLElementwiseMax::validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, const ActivationLayerInfo &act_info) { - return CLArithmeticOperationKernel::validate(ArithmeticOperation::MAX, input1, input2, output); + return CLArithmeticOperationKernel::validate(ArithmeticOperation::MAX, input1, input2, output, act_info); } -void CLElementwiseMin::configure(ICLTensor *input1, ICLTensor *input2, ICLTensor *output) +void CLElementwiseMin::configure(ICLTensor *input1, ICLTensor *input2, ICLTensor *output, const ActivationLayerInfo &act_info) { auto k = arm_compute::support::cpp14::make_unique(); - k->configure(ArithmeticOperation::MIN, input1, input2, output); + k->configure(ArithmeticOperation::MIN, input1, input2, output, act_info); _kernel = std::move(k); configure_border_handler(_border_handler, _kernel->border_size(), input1, input2, output); } -Status CLElementwiseMin::validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output) +Status CLElementwiseMin::validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, const ActivationLayerInfo &act_info) { - return CLArithmeticOperationKernel::validate(ArithmeticOperation::MIN, input1, input2, output); + return CLArithmeticOperationKernel::validate(ArithmeticOperation::MIN, input1, input2, output, act_info); } -void CLElementwiseSquaredDiff::configure(ICLTensor *input1, ICLTensor *input2, ICLTensor *output) +void CLElementwiseSquaredDiff::configure(ICLTensor *input1, ICLTensor *input2, ICLTensor *output, const ActivationLayerInfo &act_info) { auto k = arm_compute::support::cpp14::make_unique(); - k->configure(ArithmeticOperation::SQUARED_DIFF, input1, input2, output); + k->configure(ArithmeticOperation::SQUARED_DIFF, input1, input2, output, act_info); _kernel = std::move(k); configure_border_handler(_border_handler, _kernel->border_size(), input1, input2, output); } -Status CLElementwiseSquaredDiff::validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output) +Status CLElementwiseSquaredDiff::validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, const ActivationLayerInfo &act_info) { - return CLArithmeticOperationKernel::validate(ArithmeticOperation::SQUARED_DIFF, input1, input2, output); + return CLArithmeticOperationKernel::validate(ArithmeticOperation::SQUARED_DIFF, input1, input2, output, act_info); } -void CLElementwisePower::configure(ICLTensor *input1, ICLTensor *input2, ICLTensor *output) +void CLElementwisePower::configure(ICLTensor *input1, ICLTensor *input2, ICLTensor *output, const ActivationLayerInfo &act_info) { auto k = arm_compute::support::cpp14::make_unique(); - k->configure(ArithmeticOperation::POWER, input1, input2, output); + k->configure(ArithmeticOperation::POWER, input1, input2, output, act_info); _kernel = std::move(k); configure_border_handler(_border_handler, _kernel->border_size(), input1, input2, output); } -Status CLElementwisePower::validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output) +Status CLElementwisePower::validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, const ActivationLayerInfo &act_info) { - return CLArithmeticOperationKernel::validate(ArithmeticOperation::POWER, input1, input2, output); + return CLArithmeticOperationKernel::validate(ArithmeticOperation::POWER, input1, input2, output, act_info); } } // namespace arm_compute diff --git a/src/runtime/CL/functions/CLPixelWiseMultiplication.cpp b/src/runtime/CL/functions/CLPixelWiseMultiplication.cpp index c1c971816c..b527922d2b 100644 --- a/src/runtime/CL/functions/CLPixelWiseMultiplication.cpp +++ b/src/runtime/CL/functions/CLPixelWiseMultiplication.cpp @@ -32,10 +32,10 @@ namespace arm_compute { void CLPixelWiseMultiplication::configure(ICLTensor *input1, ICLTensor *input2, ICLTensor *output, float scale, - ConvertPolicy overflow_policy, RoundingPolicy rounding_policy) + ConvertPolicy overflow_policy, RoundingPolicy rounding_policy, const ActivationLayerInfo &act_info) { auto k = arm_compute::support::cpp14::make_unique(); - k->configure(input1, input2, output, scale, overflow_policy, rounding_policy); + k->configure(input1, input2, output, scale, overflow_policy, rounding_policy, act_info); _kernel = std::move(k); if(output->info()->dimension(0) > 1) @@ -50,15 +50,15 @@ void CLPixelWiseMultiplication::configure(ICLTensor *input1, ICLTensor *input2, } Status CLPixelWiseMultiplication::validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, float scale, - ConvertPolicy overflow_policy, RoundingPolicy rounding_policy) + ConvertPolicy overflow_policy, RoundingPolicy rounding_policy, const ActivationLayerInfo &act_info) { - return CLPixelWiseMultiplicationKernel::validate(input1, input2, output, scale, overflow_policy, rounding_policy); + return CLPixelWiseMultiplicationKernel::validate(input1, input2, output, scale, overflow_policy, rounding_policy, act_info); } -void CLComplexPixelWiseMultiplication::configure(ICLTensor *input1, ICLTensor *input2, ICLTensor *output) +void CLComplexPixelWiseMultiplication::configure(ICLTensor *input1, ICLTensor *input2, ICLTensor *output, const ActivationLayerInfo &act_info) { auto k = arm_compute::support::cpp14::make_unique(); - k->configure(input1, input2, output); + k->configure(input1, input2, output, act_info); _kernel = std::move(k); if(output->info()->dimension(0) > 1) @@ -72,8 +72,8 @@ void CLComplexPixelWiseMultiplication::configure(ICLTensor *input1, ICLTensor *i } } -Status CLComplexPixelWiseMultiplication::validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output) +Status CLComplexPixelWiseMultiplication::validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, const ActivationLayerInfo &act_info) { - return CLComplexPixelWiseMultiplicationKernel::validate(input1, input2, output); + return CLComplexPixelWiseMultiplicationKernel::validate(input1, input2, output, act_info); } } // namespace arm_compute diff --git a/src/runtime/GLES_COMPUTE/functions/GCArithmeticAddition.cpp b/src/runtime/GLES_COMPUTE/functions/GCArithmeticAddition.cpp index cd9c8ddad2..b0d8a3cf9f 100755 --- a/src/runtime/GLES_COMPUTE/functions/GCArithmeticAddition.cpp +++ b/src/runtime/GLES_COMPUTE/functions/GCArithmeticAddition.cpp @@ -30,14 +30,16 @@ using namespace arm_compute; -void GCArithmeticAddition::configure(const IGCTensor *input1, const IGCTensor *input2, IGCTensor *output, ConvertPolicy policy) +void GCArithmeticAddition::configure(const IGCTensor *input1, const IGCTensor *input2, IGCTensor *output, ConvertPolicy policy, const ActivationLayerInfo &act_info) { + ARM_COMPUTE_UNUSED(act_info); auto k = arm_compute::support::cpp14::make_unique(); k->configure(input1, input2, output, policy); _kernel = std::move(k); } -Status GCArithmeticAddition::validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, ConvertPolicy policy) +Status GCArithmeticAddition::validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, ConvertPolicy policy, const ActivationLayerInfo &act_info) { + ARM_COMPUTE_RETURN_ERROR_ON(act_info.enabled()); return GCArithmeticAdditionKernel::validate(input1, input2, output, policy); } diff --git a/src/runtime/GLES_COMPUTE/functions/GCPixelWiseMultiplication.cpp b/src/runtime/GLES_COMPUTE/functions/GCPixelWiseMultiplication.cpp index 126476d2a9..1075f0b5be 100755 --- a/src/runtime/GLES_COMPUTE/functions/GCPixelWiseMultiplication.cpp +++ b/src/runtime/GLES_COMPUTE/functions/GCPixelWiseMultiplication.cpp @@ -30,8 +30,9 @@ using namespace arm_compute; -void GCPixelWiseMultiplication::configure(const IGCTensor *input1, const IGCTensor *input2, IGCTensor *output, float scale) +void GCPixelWiseMultiplication::configure(const IGCTensor *input1, const IGCTensor *input2, IGCTensor *output, float scale, const ActivationLayerInfo &act_info) { + ARM_COMPUTE_UNUSED(act_info); auto k = arm_compute::support::cpp14::make_unique(); k->configure(input1, input2, output, scale); _kernel = std::move(k); diff --git a/src/runtime/NEON/functions/NEArithmeticAddition.cpp b/src/runtime/NEON/functions/NEArithmeticAddition.cpp index 6d0b207cf1..06c71db1bd 100644 --- a/src/runtime/NEON/functions/NEArithmeticAddition.cpp +++ b/src/runtime/NEON/functions/NEArithmeticAddition.cpp @@ -31,14 +31,16 @@ namespace arm_compute { -void NEArithmeticAddition::configure(ITensor *input1, ITensor *input2, ITensor *output, ConvertPolicy policy) +void NEArithmeticAddition::configure(ITensor *input1, ITensor *input2, ITensor *output, ConvertPolicy policy, const ActivationLayerInfo &act_info) { + ARM_COMPUTE_UNUSED(act_info); auto k = arm_compute::support::cpp14::make_unique(); k->configure(input1, input2, output, policy); _kernel = std::move(k); } -Status NEArithmeticAddition::validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, ConvertPolicy policy) +Status NEArithmeticAddition::validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, ConvertPolicy policy, const ActivationLayerInfo &act_info) { + ARM_COMPUTE_RETURN_ERROR_ON(act_info.enabled()); return NEArithmeticAdditionKernel::validate(input1, input2, output, policy); } } // namespace arm_compute diff --git a/src/runtime/NEON/functions/NEArithmeticSubtraction.cpp b/src/runtime/NEON/functions/NEArithmeticSubtraction.cpp index 0ad87383ce..454adc336b 100644 --- a/src/runtime/NEON/functions/NEArithmeticSubtraction.cpp +++ b/src/runtime/NEON/functions/NEArithmeticSubtraction.cpp @@ -31,8 +31,9 @@ namespace arm_compute { -void NEArithmeticSubtraction::configure(ITensor *input1, ITensor *input2, ITensor *output, ConvertPolicy policy) +void NEArithmeticSubtraction::configure(ITensor *input1, ITensor *input2, ITensor *output, ConvertPolicy policy, const ActivationLayerInfo &act_info) { + ARM_COMPUTE_UNUSED(act_info); auto k = arm_compute::support::cpp14::make_unique(); k->configure(input1, input2, output, policy); _kernel = std::move(k); @@ -48,8 +49,9 @@ void NEArithmeticSubtraction::configure(ITensor *input1, ITensor *input2, ITenso } } -Status NEArithmeticSubtraction::validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, ConvertPolicy policy) +Status NEArithmeticSubtraction::validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, ConvertPolicy policy, const ActivationLayerInfo &act_info) { + ARM_COMPUTE_RETURN_ERROR_ON(act_info.enabled()); return NEArithmeticSubtractionKernel::validate(input1, input2, output, policy); } } // namespace arm_compute diff --git a/src/runtime/NEON/functions/NEElementwiseOperators.cpp b/src/runtime/NEON/functions/NEElementwiseOperators.cpp index 0ba0ddbe3a..7451c6ff2b 100644 --- a/src/runtime/NEON/functions/NEElementwiseOperators.cpp +++ b/src/runtime/NEON/functions/NEElementwiseOperators.cpp @@ -32,15 +32,17 @@ namespace arm_compute { -void NEElementwiseMax::configure(ITensor *input1, ITensor *input2, ITensor *output) +void NEElementwiseMax::configure(ITensor *input1, ITensor *input2, ITensor *output, const ActivationLayerInfo &act_info) { + ARM_COMPUTE_UNUSED(act_info); auto k = arm_compute::support::cpp14::make_unique(); k->configure(ArithmeticOperation::MAX, input1, input2, output); _kernel = std::move(k); } -Status NEElementwiseMax::validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output) +Status NEElementwiseMax::validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, const ActivationLayerInfo &act_info) { + ARM_COMPUTE_RETURN_ERROR_ON(act_info.enabled()); if(input1->data_type() == DataType::QASYMM8_SIGNED) { ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_TYPES(input1, output); @@ -49,15 +51,17 @@ Status NEElementwiseMax::validate(const ITensorInfo *input1, const ITensorInfo * return NEArithmeticOperationKernel::validate(ArithmeticOperation::MAX, input1, input2, output); } -void NEElementwiseMin::configure(ITensor *input1, ITensor *input2, ITensor *output) +void NEElementwiseMin::configure(ITensor *input1, ITensor *input2, ITensor *output, const ActivationLayerInfo &act_info) { + ARM_COMPUTE_UNUSED(act_info); auto k = arm_compute::support::cpp14::make_unique(); k->configure(ArithmeticOperation::MIN, input1, input2, output); _kernel = std::move(k); } -Status NEElementwiseMin::validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output) +Status NEElementwiseMin::validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, const ActivationLayerInfo &act_info) { + ARM_COMPUTE_RETURN_ERROR_ON(act_info.enabled()); if(input1->data_type() == DataType::QASYMM8_SIGNED) { ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_TYPES(input1, output); @@ -66,39 +70,45 @@ Status NEElementwiseMin::validate(const ITensorInfo *input1, const ITensorInfo * return NEArithmeticOperationKernel::validate(ArithmeticOperation::MIN, input1, input2, output); } -void NEElementwiseSquaredDiff::configure(ITensor *input1, ITensor *input2, ITensor *output) +void NEElementwiseSquaredDiff::configure(ITensor *input1, ITensor *input2, ITensor *output, const ActivationLayerInfo &act_info) { + ARM_COMPUTE_UNUSED(act_info); auto k = arm_compute::support::cpp14::make_unique(); k->configure(ArithmeticOperation::SQUARED_DIFF, input1, input2, output); _kernel = std::move(k); } -Status NEElementwiseSquaredDiff::validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output) +Status NEElementwiseSquaredDiff::validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, const ActivationLayerInfo &act_info) { + ARM_COMPUTE_RETURN_ERROR_ON(act_info.enabled()); return NEArithmeticOperationKernel::validate(ArithmeticOperation::SQUARED_DIFF, input1, input2, output); } -void NEElementwiseDivision::configure(ITensor *input1, ITensor *input2, ITensor *output) +void NEElementwiseDivision::configure(ITensor *input1, ITensor *input2, ITensor *output, const ActivationLayerInfo &act_info) { + ARM_COMPUTE_UNUSED(act_info); auto k = arm_compute::support::cpp14::make_unique(); k->configure(input1, input2, output); _kernel = std::move(k); } -Status NEElementwiseDivision::validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output) +Status NEElementwiseDivision::validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, const ActivationLayerInfo &act_info) { + ARM_COMPUTE_RETURN_ERROR_ON(act_info.enabled()); return NEDivisionOperationKernel::validate(input1, input2, output); } -void NEElementwisePower::configure(ITensor *input1, ITensor *input2, ITensor *output) +void NEElementwisePower::configure(ITensor *input1, ITensor *input2, ITensor *output, const ActivationLayerInfo &act_info) { + ARM_COMPUTE_UNUSED(act_info); auto k = arm_compute::support::cpp14::make_unique(); k->configure(input1, input2, output); _kernel = std::move(k); } -Status NEElementwisePower::validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output) +Status NEElementwisePower::validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, const ActivationLayerInfo &act_info) { + ARM_COMPUTE_RETURN_ERROR_ON(act_info.enabled()); return NEPowerOperationKernel::validate(input1, input2, output); } diff --git a/src/runtime/NEON/functions/NEPixelWiseMultiplication.cpp b/src/runtime/NEON/functions/NEPixelWiseMultiplication.cpp index e2516e420c..eaf233b9ed 100644 --- a/src/runtime/NEON/functions/NEPixelWiseMultiplication.cpp +++ b/src/runtime/NEON/functions/NEPixelWiseMultiplication.cpp @@ -31,8 +31,10 @@ namespace arm_compute { -void NEPixelWiseMultiplication::configure(ITensor *input1, ITensor *input2, ITensor *output, float scale, ConvertPolicy overflow_policy, RoundingPolicy rounding_policy) +void NEPixelWiseMultiplication::configure(ITensor *input1, ITensor *input2, ITensor *output, float scale, ConvertPolicy overflow_policy, RoundingPolicy rounding_policy, + const ActivationLayerInfo &act_info) { + ARM_COMPUTE_UNUSED(act_info); auto k = arm_compute::support::cpp14::make_unique(); k->configure(input1, input2, output, scale, overflow_policy, rounding_policy); _kernel = std::move(k); @@ -47,13 +49,16 @@ void NEPixelWiseMultiplication::configure(ITensor *input1, ITensor *input2, ITen } } } -Status NEPixelWiseMultiplication::validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, float scale, ConvertPolicy overflow_policy, RoundingPolicy rounding_policy) +Status NEPixelWiseMultiplication::validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, float scale, ConvertPolicy overflow_policy, RoundingPolicy rounding_policy, + const ActivationLayerInfo &act_info) { + ARM_COMPUTE_RETURN_ERROR_ON(act_info.enabled()); return NEPixelWiseMultiplicationKernel::validate(input1, input2, output, scale, overflow_policy, rounding_policy); } -void NEComplexPixelWiseMultiplication::configure(ITensor *input1, ITensor *input2, ITensor *output) +void NEComplexPixelWiseMultiplication::configure(ITensor *input1, ITensor *input2, ITensor *output, const ActivationLayerInfo &act_info) { + ARM_COMPUTE_UNUSED(act_info); auto k = arm_compute::support::cpp14::make_unique(); k->configure(input1, input2, output); _kernel = std::move(k); @@ -69,8 +74,9 @@ void NEComplexPixelWiseMultiplication::configure(ITensor *input1, ITensor *input } } -Status NEComplexPixelWiseMultiplication::validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output) +Status NEComplexPixelWiseMultiplication::validate(const ITensorInfo *input1, const ITensorInfo *input2, const ITensorInfo *output, const ActivationLayerInfo &act_info) { + ARM_COMPUTE_RETURN_ERROR_ON(act_info.enabled()); return NEComplexPixelWiseMultiplicationKernel::validate(input1, input2, output); } diff --git a/tests/datasets/ShapeDatasets.h b/tests/datasets/ShapeDatasets.h index 2a2047480f..087342d3b8 100644 --- a/tests/datasets/ShapeDatasets.h +++ b/tests/datasets/ShapeDatasets.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2019 ARM Limited. + * Copyright (c) 2017-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -156,7 +156,7 @@ public: } }; -/** Data set containing small tensor shapes. */ +/** Data set containing tiny tensor shapes. */ class TinyShapes final : public ShapeDataset { public: @@ -190,6 +190,25 @@ public: } }; +/** Data set containing pairs of tiny tensor shapes that are broadcast compatible. */ +class TinyShapesBroadcast final : public framework::dataset::ZipDataset +{ +public: + TinyShapesBroadcast() + : ZipDataset( + ShapeDataset("Shape0", + { + TensorShape{ 9U, 9U }, + TensorShape{ 10U, 2U, 14U, 2U }, + }), + ShapeDataset("Shape1", + { + TensorShape{ 9U, 1U, 9U }, + TensorShape{ 10U }, + })) + { + } +}; /** Data set containing pairs of small tensor shapes that are broadcast compatible. */ class SmallShapesBroadcast final : public framework::dataset::ZipDataset { diff --git a/tests/validation/CL/ArithmeticAddition.cpp b/tests/validation/CL/ArithmeticAddition.cpp index 5702360e45..41415ee481 100644 --- a/tests/validation/CL/ArithmeticAddition.cpp +++ b/tests/validation/CL/ArithmeticAddition.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2019 ARM Limited. + * Copyright (c) 2017-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -51,8 +51,8 @@ const auto ArithmeticAdditionQASYMM8Dataset = combine(combine(framework::dataset framework::dataset::make("DataType", DataType::QASYMM8)); const auto ArithmeticAdditionQASYMM8SignedDataset = combine(combine(framework::dataset::make("DataType", DataType::QASYMM8_SIGNED), framework::dataset::make("DataType", DataType::QASYMM8_SIGNED)), - framework::dataset::make("DataType", - DataType::QASYMM8_SIGNED)); + framework::dataset::make("DataType", + DataType::QASYMM8_SIGNED)); const auto ArithmeticAdditionQSYMM16Dataset = combine(combine(framework::dataset::make("DataType", DataType::QSYMM16), framework::dataset::make("DataType", DataType::QSYMM16)), framework::dataset::make("DataType", DataType::QSYMM16)); @@ -62,6 +62,13 @@ const auto ArithmeticAdditionFP16Dataset = combine(combine(framework::dataset::m framework::dataset::make("DataType", DataType::F16)); const auto ArithmeticAdditionFP32Dataset = combine(combine(framework::dataset::make("DataType", DataType::F32), framework::dataset::make("DataType", DataType::F32)), framework::dataset::make("DataType", DataType::F32)); +const auto EmptyActivationFunctionsDataset = framework::dataset::make("ActivationInfo", +{ ActivationLayerInfo() }); +const auto ActivationFunctionsDataset = framework::dataset::make("ActivationInfo", +{ + ActivationLayerInfo(ActivationLayerInfo::ActivationFunction::BOUNDED_RELU, 0.75f, 0.25f), + ActivationLayerInfo(ActivationLayerInfo::ActivationFunction::LOGISTIC, 0.75f, 0.25f) +}); } // namespace TEST_SUITE(CL) @@ -284,10 +291,21 @@ FIXTURE_DATA_TEST_CASE(RunSmall, CLArithmeticAdditionQuantizedFixture, TEST_SUITE_END() // QSYMM16 TEST_SUITE_END() // Quantized +template +using CLArithmeticAdditionFloatFixture = ArithmeticAdditionValidationFloatFixture; + TEST_SUITE(Float) TEST_SUITE(FP16) -FIXTURE_DATA_TEST_CASE(RunSmall, CLArithmeticAdditionFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapes(), ArithmeticAdditionFP16Dataset), - framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP }))) +FIXTURE_DATA_TEST_CASE(RunSmall, CLArithmeticAdditionFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(datasets::SmallShapes(), ArithmeticAdditionFP16Dataset), + framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP })), + EmptyActivationFunctionsDataset)) +{ + // Validate output + validate(CLAccessor(_target), _reference); +} +FIXTURE_DATA_TEST_CASE(RunWithActivation, CLArithmeticAdditionFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(datasets::TinyShapes(), ArithmeticAdditionFP16Dataset), + framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP })), + ActivationFunctionsDataset)) { // Validate output validate(CLAccessor(_target), _reference); @@ -318,34 +336,53 @@ DATA_TEST_CASE(Configuration, framework::DatasetMode::ALL, combine(datasets::Sma validate(dst.info()->padding(), padding); } -FIXTURE_DATA_TEST_CASE(RunSmall, CLArithmeticAdditionFixture, framework::DatasetMode::PRECOMMIT, combine(combine(datasets::SmallShapes(), ArithmeticAdditionFP32Dataset), - framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP }))) +FIXTURE_DATA_TEST_CASE(RunSmall, CLArithmeticAdditionFloatFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(datasets::SmallShapes(), ArithmeticAdditionFP32Dataset), + framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP })), + EmptyActivationFunctionsDataset)) +{ + // Validate output + validate(CLAccessor(_target), _reference); +} +FIXTURE_DATA_TEST_CASE(RunWithActivation, CLArithmeticAdditionFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(datasets::TinyShapes(), ArithmeticAdditionFP32Dataset), + framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP })), + ActivationFunctionsDataset)) { // Validate output validate(CLAccessor(_target), _reference); } -FIXTURE_DATA_TEST_CASE(RunLarge, CLArithmeticAdditionFixture, framework::DatasetMode::NIGHTLY, combine(combine(datasets::LargeShapes(), ArithmeticAdditionFP32Dataset), - framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP }))) +FIXTURE_DATA_TEST_CASE(RunLarge, CLArithmeticAdditionFloatFixture, framework::DatasetMode::NIGHTLY, combine(combine(combine(datasets::LargeShapes(), ArithmeticAdditionFP32Dataset), + framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP })), + EmptyActivationFunctionsDataset)) { // Validate output validate(CLAccessor(_target), _reference); } template -using CLArithmeticAdditionBroadcastFixture = ArithmeticAdditionBroadcastValidationFixture; +using CLArithmeticAdditionBroadcastFloatFixture = ArithmeticAdditionBroadcastValidationFloatFixture; -FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, CLArithmeticAdditionBroadcastFixture, framework::DatasetMode::PRECOMMIT, combine(combine(datasets::SmallShapesBroadcast(), +FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, CLArithmeticAdditionBroadcastFloatFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(datasets::SmallShapesBroadcast(), + ArithmeticAdditionFP32Dataset), + framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP })), + EmptyActivationFunctionsDataset)) +{ + // Validate output + validate(CLAccessor(_target), _reference); +} +FIXTURE_DATA_TEST_CASE(RunWithActivationBroadcast, CLArithmeticAdditionBroadcastFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(datasets::TinyShapesBroadcast(), ArithmeticAdditionFP32Dataset), - framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP }))) + framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP })), + ActivationFunctionsDataset)) { // Validate output validate(CLAccessor(_target), _reference); } -FIXTURE_DATA_TEST_CASE(RunLargeBroadcast, CLArithmeticAdditionBroadcastFixture, framework::DatasetMode::NIGHTLY, combine(combine(datasets::LargeShapesBroadcast(), +FIXTURE_DATA_TEST_CASE(RunLargeBroadcast, CLArithmeticAdditionBroadcastFloatFixture, framework::DatasetMode::NIGHTLY, combine(combine(combine(datasets::LargeShapesBroadcast(), ArithmeticAdditionFP32Dataset), - framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP }))) + framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP })), + EmptyActivationFunctionsDataset)) { // Validate output validate(CLAccessor(_target), _reference); diff --git a/tests/validation/CL/ArithmeticDivision.cpp b/tests/validation/CL/ArithmeticDivision.cpp index 8df770ab98..d970c31daa 100644 --- a/tests/validation/CL/ArithmeticDivision.cpp +++ b/tests/validation/CL/ArithmeticDivision.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2019 ARM Limited. + * Copyright (c) 2018-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -52,6 +52,13 @@ const auto ArithmeticDivisionFP16Dataset = combine(combine(framework::dataset::m framework::dataset::make("DataType", DataType::F16)); const auto ArithmeticDivisionFP32Dataset = combine(combine(framework::dataset::make("DataType", DataType::F32), framework::dataset::make("DataType", DataType::F32)), framework::dataset::make("DataType", DataType::F32)); +const auto EmptyActivationFunctionsDataset = framework::dataset::make("ActivationInfo", +{ ActivationLayerInfo() }); +const auto ActivationFunctionsDataset = framework::dataset::make("ActivationInfo", +{ + ActivationLayerInfo(ActivationLayerInfo::ActivationFunction::BOUNDED_RELU, 0.75f, 0.25f), + ActivationLayerInfo(ActivationLayerInfo::ActivationFunction::LOGISTIC, 0.75f, 0.25f) +}); } // namespace TEST_SUITE(CL) @@ -87,11 +94,18 @@ DATA_TEST_CASE(Validate, framework::DatasetMode::ALL, zip(zip(zip( // *INDENT-ON* template -using CLArithmeticDivisionFixture = ArithmeticDivisionValidationFixture; +using CLArithmeticDivisionFloatFixture = ArithmeticDivisionValidationFloatFixture; TEST_SUITE(Float) TEST_SUITE(FP16) -FIXTURE_DATA_TEST_CASE(RunSmall, CLArithmeticDivisionFixture, framework::DatasetMode::ALL, combine(datasets::SmallShapes(), ArithmeticDivisionFP16Dataset)) +FIXTURE_DATA_TEST_CASE(RunSmall, CLArithmeticDivisionFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapes(), ArithmeticDivisionFP16Dataset), + EmptyActivationFunctionsDataset)) +{ + // Validate output + validate(CLAccessor(_target), _reference, tolerance_fp16, 0.01); +} +FIXTURE_DATA_TEST_CASE(RunWithActivation, CLArithmeticDivisionFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::TinyShapes(), ArithmeticDivisionFP16Dataset), + ActivationFunctionsDataset)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp16, 0.01); @@ -122,30 +136,47 @@ DATA_TEST_CASE(Configuration, framework::DatasetMode::ALL, datasets::SmallShapes validate(dst.info()->padding(), padding); } -FIXTURE_DATA_TEST_CASE(RunSmall, CLArithmeticDivisionFixture, framework::DatasetMode::PRECOMMIT, combine(datasets::SmallShapes(), ArithmeticDivisionFP32Dataset)) +FIXTURE_DATA_TEST_CASE(RunSmall, CLArithmeticDivisionFloatFixture, framework::DatasetMode::PRECOMMIT, combine(combine(datasets::SmallShapes(), ArithmeticDivisionFP32Dataset), + EmptyActivationFunctionsDataset)) +{ + // Validate output + validate(CLAccessor(_target), _reference, tolerance_fp32); +} +FIXTURE_DATA_TEST_CASE(RunWithActivation, CLArithmeticDivisionFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::TinyShapes(), ArithmeticDivisionFP32Dataset), + ActivationFunctionsDataset)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp32); } -FIXTURE_DATA_TEST_CASE(RunLarge, CLArithmeticDivisionFixture, framework::DatasetMode::NIGHTLY, combine(datasets::LargeShapes(), ArithmeticDivisionFP32Dataset)) +FIXTURE_DATA_TEST_CASE(RunLarge, CLArithmeticDivisionFloatFixture, framework::DatasetMode::NIGHTLY, combine(combine(datasets::LargeShapes(), ArithmeticDivisionFP32Dataset), + EmptyActivationFunctionsDataset)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp32); } template -using CLArithmeticDivisionBroadcastFixture = ArithmeticDivisionBroadcastValidationFixture; +using CLArithmeticDivisionBroadcastFloatFixture = ArithmeticDivisionBroadcastValidationFloatFixture; -FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, CLArithmeticDivisionBroadcastFixture, framework::DatasetMode::PRECOMMIT, combine(datasets::SmallShapesBroadcast(), - ArithmeticDivisionFP32Dataset)) +FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, CLArithmeticDivisionBroadcastFloatFixture, framework::DatasetMode::PRECOMMIT, combine(combine(datasets::SmallShapesBroadcast(), + ArithmeticDivisionFP32Dataset), + EmptyActivationFunctionsDataset)) +{ + // Validate output + validate(CLAccessor(_target), _reference, tolerance_fp32); +} +FIXTURE_DATA_TEST_CASE(RunWithActivationBroadcast, CLArithmeticDivisionBroadcastFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::TinyShapesBroadcast(), + ArithmeticDivisionFP32Dataset), + ActivationFunctionsDataset)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp32); } -FIXTURE_DATA_TEST_CASE(RunLargeBroadcast, CLArithmeticDivisionBroadcastFixture, framework::DatasetMode::NIGHTLY, combine(datasets::LargeShapesBroadcast(), - ArithmeticDivisionFP32Dataset)) +FIXTURE_DATA_TEST_CASE(RunLargeBroadcast, CLArithmeticDivisionBroadcastFloatFixture, framework::DatasetMode::NIGHTLY, combine(combine(datasets::LargeShapesBroadcast(), + ArithmeticDivisionFP32Dataset), + EmptyActivationFunctionsDataset)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp32); diff --git a/tests/validation/CL/ArithmeticSubtraction.cpp b/tests/validation/CL/ArithmeticSubtraction.cpp index 592a7ed1a4..897ae1ab09 100644 --- a/tests/validation/CL/ArithmeticSubtraction.cpp +++ b/tests/validation/CL/ArithmeticSubtraction.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2019 ARM Limited. + * Copyright (c) 2017-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -52,8 +52,8 @@ const auto ArithmeticSubtractionQASYMM8Dataset = combine(combine(framework::data framework::dataset::make("DataType", DataType::QASYMM8)); const auto ArithmeticSubtractionQASYMM8SignedDataset = combine(combine(framework::dataset::make("DataType", DataType::QASYMM8_SIGNED), framework::dataset::make("DataType", DataType::QASYMM8_SIGNED)), - framework::dataset::make("DataType", - DataType::QASYMM8_SIGNED)); + framework::dataset::make("DataType", + DataType::QASYMM8_SIGNED)); const auto ArithmeticSubtractionQSYMM16Dataset = combine(combine(framework::dataset::make("DataType", DataType::QSYMM16), framework::dataset::make("DataType", DataType::QSYMM16)), framework::dataset::make("DataType", DataType::QSYMM16)); @@ -63,6 +63,13 @@ const auto ArithmeticSubtractionFP16Dataset = combine(combine(framework::dataset framework::dataset::make("DataType", DataType::F16)); const auto ArithmeticSubtractionFP32Dataset = combine(combine(framework::dataset::make("DataType", DataType::F32), framework::dataset::make("DataType", DataType::F32)), framework::dataset::make("DataType", DataType::F32)); +const auto EmptyActivationFunctionsDataset = framework::dataset::make("ActivationInfo", +{ ActivationLayerInfo() }); +const auto ActivationFunctionsDataset = framework::dataset::make("ActivationInfo", +{ + ActivationLayerInfo(ActivationLayerInfo::ActivationFunction::BOUNDED_RELU, 0.75f, 0.25f), + ActivationLayerInfo(ActivationLayerInfo::ActivationFunction::LOGISTIC, 0.75f, 0.25f) +}); } // namespace TEST_SUITE(CL) @@ -285,10 +292,21 @@ FIXTURE_DATA_TEST_CASE(RunSmall, CLArithmeticSubtractionQuantizedFixture +using CLArithmeticSubtractionFloatFixture = ArithmeticSubtractionValidationFloatFixture; + TEST_SUITE(Float) TEST_SUITE(FP16) -FIXTURE_DATA_TEST_CASE(RunSmall, CLArithmeticSubtractionFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapes(), ArithmeticSubtractionFP16Dataset), - framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP }))) +FIXTURE_DATA_TEST_CASE(RunSmall, CLArithmeticSubtractionFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(datasets::SmallShapes(), ArithmeticSubtractionFP16Dataset), + framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP })), + EmptyActivationFunctionsDataset)) +{ + // Validate output + validate(CLAccessor(_target), _reference); +} +FIXTURE_DATA_TEST_CASE(RunWithActivation, CLArithmeticSubtractionFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(datasets::TinyShapes(), ArithmeticSubtractionFP16Dataset), + framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP })), + ActivationFunctionsDataset)) { // Validate output validate(CLAccessor(_target), _reference); @@ -319,34 +337,53 @@ DATA_TEST_CASE(Configuration, framework::DatasetMode::ALL, combine(datasets::Sma validate(dst.info()->padding(), padding); } -FIXTURE_DATA_TEST_CASE(RunSmall, CLArithmeticSubtractionFixture, framework::DatasetMode::PRECOMMIT, combine(combine(datasets::SmallShapes(), ArithmeticSubtractionFP32Dataset), - framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP }))) +FIXTURE_DATA_TEST_CASE(RunSmall, CLArithmeticSubtractionFloatFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(datasets::SmallShapes(), ArithmeticSubtractionFP32Dataset), + framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP })), + EmptyActivationFunctionsDataset)) +{ + // Validate output + validate(CLAccessor(_target), _reference); +} +FIXTURE_DATA_TEST_CASE(RunWithActivation, CLArithmeticSubtractionFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(datasets::TinyShapes(), ArithmeticSubtractionFP32Dataset), + framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP })), + ActivationFunctionsDataset)) { // Validate output validate(CLAccessor(_target), _reference); } -FIXTURE_DATA_TEST_CASE(RunLarge, CLArithmeticSubtractionFixture, framework::DatasetMode::NIGHTLY, combine(combine(datasets::LargeShapes(), ArithmeticSubtractionFP32Dataset), - framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP }))) +FIXTURE_DATA_TEST_CASE(RunLarge, CLArithmeticSubtractionFloatFixture, framework::DatasetMode::NIGHTLY, combine(combine(combine(datasets::LargeShapes(), ArithmeticSubtractionFP32Dataset), + framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP })), + EmptyActivationFunctionsDataset)) { // Validate output validate(CLAccessor(_target), _reference); } template -using CLArithmeticSubtractionBroadcastFixture = ArithmeticSubtractionBroadcastValidationFixture; +using CLArithmeticSubtractionBroadcastFloatFixture = ArithmeticSubtractionBroadcastValidationFloatFixture; -FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, CLArithmeticSubtractionBroadcastFixture, framework::DatasetMode::PRECOMMIT, combine(combine(datasets::SmallShapesBroadcast(), +FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, CLArithmeticSubtractionBroadcastFloatFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(datasets::SmallShapesBroadcast(), + ArithmeticSubtractionFP32Dataset), + framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP })), + EmptyActivationFunctionsDataset)) +{ + // Validate output + validate(CLAccessor(_target), _reference); +} +FIXTURE_DATA_TEST_CASE(RunWithActivationBroadcast, CLArithmeticSubtractionBroadcastFloatFixture, framework::DatasetMode::ALL, combine(combine(combine(datasets::TinyShapesBroadcast(), ArithmeticSubtractionFP32Dataset), - framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP }))) + framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP })), + ActivationFunctionsDataset)) { // Validate output validate(CLAccessor(_target), _reference); } -FIXTURE_DATA_TEST_CASE(RunLargeBroadcast, CLArithmeticSubtractionBroadcastFixture, framework::DatasetMode::NIGHTLY, combine(combine(datasets::LargeShapesBroadcast(), +FIXTURE_DATA_TEST_CASE(RunLargeBroadcast, CLArithmeticSubtractionBroadcastFloatFixture, framework::DatasetMode::NIGHTLY, combine(combine(combine(datasets::LargeShapesBroadcast(), ArithmeticSubtractionFP32Dataset), - framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP }))) + framework::dataset::make("ConvertPolicy", { ConvertPolicy::SATURATE, ConvertPolicy::WRAP })), + EmptyActivationFunctionsDataset)) { // Validate output validate(CLAccessor(_target), _reference); diff --git a/tests/validation/CL/ElementwiseMax.cpp b/tests/validation/CL/ElementwiseMax.cpp index c7b4284e19..879e732cb0 100644 --- a/tests/validation/CL/ElementwiseMax.cpp +++ b/tests/validation/CL/ElementwiseMax.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2019 ARM Limited. + * Copyright (c) 2018-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -54,8 +54,8 @@ const auto ElementwiseMaxQASYMM8Dataset = combine(combine(framework::dataset::ma framework::dataset::make("DataType", DataType::QASYMM8)); const auto ElementwiseMaxQASYMM8SignedDataset = combine(combine(framework::dataset::make("DataType", DataType::QASYMM8_SIGNED), framework::dataset::make("DataType", DataType::QASYMM8_SIGNED)), - framework::dataset::make("DataType", - DataType::QASYMM8_SIGNED)); + framework::dataset::make("DataType", + DataType::QASYMM8_SIGNED)); const auto ElementwiseMaxQSYMM16Dataset = combine(combine(framework::dataset::make("DataType", DataType::QSYMM16), framework::dataset::make("DataType", DataType::QSYMM16)), framework::dataset::make("DataType", DataType::QSYMM16)); @@ -65,6 +65,13 @@ const auto ElementwiseMaxFP16Dataset = combine(combine(framework::dataset::make( framework::dataset::make("DataType", DataType::F16)); const auto ElementwiseMaxFP32Dataset = combine(combine(framework::dataset::make("DataType", DataType::F32), framework::dataset::make("DataType", DataType::F32)), framework::dataset::make("DataType", DataType::F32)); +const auto EmptyActivationFunctionsDataset = framework::dataset::make("ActivationInfo", +{ ActivationLayerInfo() }); +const auto ActivationFunctionsDataset = framework::dataset::make("ActivationInfo", +{ + ActivationLayerInfo(ActivationLayerInfo::ActivationFunction::BOUNDED_RELU, 0.75f, 0.25f), + ActivationLayerInfo(ActivationLayerInfo::ActivationFunction::LOGISTIC, 0.75f, 0.25f) +}); } // namespace TEST_SUITE(CL) @@ -229,10 +236,10 @@ DATA_TEST_CASE(Configuration, framework::DatasetMode::ALL, datasets::SmallShapes } FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseMaxQuantizedFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(datasets::SmallShapes(), - ElementwiseMaxQASYMM8SignedDataset), - framework::dataset::make("Src0QInfo", { QuantizationInfo(5.f / 255.f, 20) })), - framework::dataset::make("Src1QInfo", { QuantizationInfo(2.f / 255.f, 10) })), - framework::dataset::make("OutQInfo", { QuantizationInfo(1.f / 255.f, 5) }))) + ElementwiseMaxQASYMM8SignedDataset), + framework::dataset::make("Src0QInfo", { QuantizationInfo(5.f / 255.f, 20) })), + framework::dataset::make("Src1QInfo", { QuantizationInfo(2.f / 255.f, 10) })), + framework::dataset::make("OutQInfo", { QuantizationInfo(1.f / 255.f, 5) }))) { // Validate output validate(CLAccessor(_target), _reference); @@ -274,9 +281,18 @@ FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseMaxQuantizedFixture, fram TEST_SUITE_END() TEST_SUITE_END() +template +using CLElementwiseMaxFloatFixture = ElementwiseMaxValidationFloatFixture; + TEST_SUITE(Float) TEST_SUITE(FP16) -FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseMaxFixture, framework::DatasetMode::ALL, combine(datasets::SmallShapes(), ElementwiseMaxFP16Dataset)) +FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseMaxFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapes(), ElementwiseMaxFP16Dataset), EmptyActivationFunctionsDataset)) +{ + // Validate output + validate(CLAccessor(_target), _reference, tolerance_fp16, 0.01); +} +FIXTURE_DATA_TEST_CASE(RunWithActivation, CLElementwiseMaxFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::TinyShapes(), ElementwiseMaxFP16Dataset), + ActivationFunctionsDataset)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp16, 0.01); @@ -307,17 +323,32 @@ DATA_TEST_CASE(Configuration, framework::DatasetMode::ALL, datasets::SmallShapes validate(dst.info()->padding(), padding); } -FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseMaxFixture, framework::DatasetMode::ALL, combine(datasets::SmallShapes(), ElementwiseMaxFP32Dataset)) +FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseMaxFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapes(), ElementwiseMaxFP32Dataset), + EmptyActivationFunctionsDataset)) +{ + // Validate output + validate(CLAccessor(_target), _reference, tolerance_fp32); +} +FIXTURE_DATA_TEST_CASE(RunWithActivation, CLElementwiseMaxFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::TinyShapes(), ElementwiseMaxFP32Dataset), + ActivationFunctionsDataset)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp32); } template -using CLElementwiseMaxBroadcastFixture = ElementwiseMaxBroadcastValidationFixture; +using CLElementwiseMaxBroadcastFloatFixture = ElementwiseMaxBroadcastValidationFloatFixture; -FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, CLElementwiseMaxBroadcastFixture, framework::DatasetMode::ALL, combine(datasets::SmallShapesBroadcast(), - ElementwiseMaxFP32Dataset)) +FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, CLElementwiseMaxBroadcastFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapesBroadcast(), + ElementwiseMaxFP32Dataset), + EmptyActivationFunctionsDataset)) +{ + // Validate output + validate(CLAccessor(_target), _reference, tolerance_fp32); +} +FIXTURE_DATA_TEST_CASE(RunWithActivationBroadcast, CLElementwiseMaxBroadcastFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::TinyShapesBroadcast(), + ElementwiseMaxFP32Dataset), + ActivationFunctionsDataset)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp32); diff --git a/tests/validation/CL/ElementwiseMin.cpp b/tests/validation/CL/ElementwiseMin.cpp index 3d6bde1c98..332fa80d72 100644 --- a/tests/validation/CL/ElementwiseMin.cpp +++ b/tests/validation/CL/ElementwiseMin.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2019 ARM Limited. + * Copyright (c) 2018-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -54,8 +54,8 @@ const auto ElementwiseMinQASYMM8Dataset = combine(combine(framework::dataset::ma framework::dataset::make("DataType", DataType::QASYMM8)); const auto ElementwiseMinQASYMM8SignedDataset = combine(combine(framework::dataset::make("DataType", DataType::QASYMM8_SIGNED), framework::dataset::make("DataType", DataType::QASYMM8_SIGNED)), - framework::dataset::make("DataType", - DataType::QASYMM8_SIGNED)); + framework::dataset::make("DataType", + DataType::QASYMM8_SIGNED)); const auto ElementwiseMinQSYMM16Dataset = combine(combine(framework::dataset::make("DataType", DataType::QSYMM16), framework::dataset::make("DataType", DataType::QSYMM16)), framework::dataset::make("DataType", DataType::QSYMM16)); @@ -65,6 +65,13 @@ const auto ElementwiseMinFP16Dataset = combine(combine(framework::dataset::make( framework::dataset::make("DataType", DataType::F16)); const auto ElementwiseMinFP32Dataset = combine(combine(framework::dataset::make("DataType", DataType::F32), framework::dataset::make("DataType", DataType::F32)), framework::dataset::make("DataType", DataType::F32)); +const auto EmptyActivationFunctionsDataset = framework::dataset::make("ActivationInfo", +{ ActivationLayerInfo() }); +const auto ActivationFunctionsDataset = framework::dataset::make("ActivationInfo", +{ + ActivationLayerInfo(ActivationLayerInfo::ActivationFunction::BOUNDED_RELU, 0.75f, 0.25f), + ActivationLayerInfo(ActivationLayerInfo::ActivationFunction::LOGISTIC, 0.75f, 0.25f) +}); } // namespace TEST_SUITE(CL) @@ -229,10 +236,10 @@ DATA_TEST_CASE(Configuration, framework::DatasetMode::ALL, datasets::SmallShapes } FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseMinQuantizedFixture, framework::DatasetMode::PRECOMMIT, combine(combine(combine(combine(datasets::SmallShapes(), - ElementwiseMinQASYMM8SignedDataset), - framework::dataset::make("Src0QInfo", { QuantizationInfo(5.f / 255.f, 20) })), - framework::dataset::make("Src1QInfo", { QuantizationInfo(2.f / 255.f, 10) })), - framework::dataset::make("OutQInfo", { QuantizationInfo(1.f / 255.f, 5) }))) + ElementwiseMinQASYMM8SignedDataset), + framework::dataset::make("Src0QInfo", { QuantizationInfo(5.f / 255.f, 20) })), + framework::dataset::make("Src1QInfo", { QuantizationInfo(2.f / 255.f, 10) })), + framework::dataset::make("OutQInfo", { QuantizationInfo(1.f / 255.f, 5) }))) { // Validate output validate(CLAccessor(_target), _reference); @@ -274,9 +281,18 @@ FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseMinQuantizedFixture, fram TEST_SUITE_END() TEST_SUITE_END() +template +using CLElementwiseMinFloatFixture = ElementwiseMinValidationFloatFixture; + TEST_SUITE(Float) TEST_SUITE(FP16) -FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseMinFixture, framework::DatasetMode::ALL, combine(datasets::SmallShapes(), ElementwiseMinFP16Dataset)) +FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseMinFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapes(), ElementwiseMinFP16Dataset), EmptyActivationFunctionsDataset)) +{ + // Validate output + validate(CLAccessor(_target), _reference, tolerance_fp16, 0.01); +} +FIXTURE_DATA_TEST_CASE(RunWithActivation, CLElementwiseMinFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::TinyShapes(), ElementwiseMinFP16Dataset), + ActivationFunctionsDataset)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp16, 0.01); @@ -307,16 +323,31 @@ DATA_TEST_CASE(Configuration, framework::DatasetMode::ALL, datasets::SmallShapes validate(dst.info()->padding(), padding); } -FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseMinFixture, framework::DatasetMode::ALL, combine(datasets::SmallShapes(), ElementwiseMinFP32Dataset)) +FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseMinFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapes(), ElementwiseMinFP32Dataset), + EmptyActivationFunctionsDataset)) +{ + // Validate output + validate(CLAccessor(_target), _reference, tolerance_fp32); +} +FIXTURE_DATA_TEST_CASE(RunWithActivation, CLElementwiseMinFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::TinyShapes(), ElementwiseMinFP32Dataset), + ActivationFunctionsDataset)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp32); } template -using CLElementwiseMinBroadcastFixture = ElementwiseMinBroadcastValidationFixture; +using CLElementwiseMinBroadcastFloatFixture = ElementwiseMinBroadcastValidationFloatFixture; -FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, CLElementwiseMinBroadcastFixture, framework::DatasetMode::ALL, combine(datasets::SmallShapesBroadcast(), - ElementwiseMinFP32Dataset)) +FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, CLElementwiseMinBroadcastFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapesBroadcast(), + ElementwiseMinFP32Dataset), + EmptyActivationFunctionsDataset)) +{ + // Validate output + validate(CLAccessor(_target), _reference, tolerance_fp32); +} +FIXTURE_DATA_TEST_CASE(RunWithActivationBroadcast, CLElementwiseMinBroadcastFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::TinyShapesBroadcast(), + ElementwiseMinFP32Dataset), + ActivationFunctionsDataset)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp32); diff --git a/tests/validation/CL/ElementwisePower.cpp b/tests/validation/CL/ElementwisePower.cpp index 46509a2870..ce4fc80bb0 100644 --- a/tests/validation/CL/ElementwisePower.cpp +++ b/tests/validation/CL/ElementwisePower.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019 ARM Limited. + * Copyright (c) 2019-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -50,6 +50,13 @@ const auto ElementwisePowerFP16Dataset = combine(combine(framework::dataset::mak framework::dataset::make("DataType", DataType::F16)); const auto ElementwisePowerFP32Dataset = combine(combine(framework::dataset::make("DataType", DataType::F32), framework::dataset::make("DataType", DataType::F32)), framework::dataset::make("DataType", DataType::F32)); +const auto EmptyActivationFunctionsDataset = framework::dataset::make("ActivationInfo", +{ ActivationLayerInfo() }); +const auto ActivationFunctionsDataset = framework::dataset::make("ActivationInfo", +{ + ActivationLayerInfo(ActivationLayerInfo::ActivationFunction::BOUNDED_RELU, 0.75f, 0.25f), + ActivationLayerInfo(ActivationLayerInfo::ActivationFunction::LOGISTIC, 0.75f, 0.25f) +}); } // namespace TEST_SUITE(CL) @@ -85,21 +92,36 @@ DATA_TEST_CASE(Validate, framework::DatasetMode::ALL, zip(zip(zip( // *INDENT-ON* template -using CLElementwisePowerFixture = ElementwisePowerValidationFixture; +using CLElementwisePowerFloatFixture = ElementwisePowerValidationFloatFixture; template -using CLElementwisePowerBroadcastFixture = ElementwisePowerBroadcastValidationFixture; +using CLElementwisePowerBroadcastFloatFixture = ElementwisePowerBroadcastValidationFloatFixture; TEST_SUITE(Float) TEST_SUITE(FP16) -FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwisePowerFixture, framework::DatasetMode::ALL, combine(datasets::SmallShapes(), ElementwisePowerFP16Dataset)) +FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwisePowerFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapes(), ElementwisePowerFP16Dataset), + EmptyActivationFunctionsDataset)) +{ + // Validate output + validate(CLAccessor(_target), _reference, tolerance_fp16, 0.01); +} +FIXTURE_DATA_TEST_CASE(RunWithActivation, CLElementwisePowerFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::TinyShapes(), ElementwisePowerFP16Dataset), + ActivationFunctionsDataset)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp16, 0.01); } -FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, CLElementwisePowerBroadcastFixture, framework::DatasetMode::ALL, combine(datasets::SmallShapesBroadcast(), - ElementwisePowerFP16Dataset)) +FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, CLElementwisePowerBroadcastFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapesBroadcast(), + ElementwisePowerFP16Dataset), + EmptyActivationFunctionsDataset)) +{ + // Validate output + validate(CLAccessor(_target), _reference, tolerance_fp16, 0.01); +} +FIXTURE_DATA_TEST_CASE(RunWithActivationBroadcast, CLElementwisePowerBroadcastFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::TinyShapesBroadcast(), + ElementwisePowerFP16Dataset), + ActivationFunctionsDataset)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp16, 0.01); @@ -107,14 +129,29 @@ FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, CLElementwisePowerBroadcastFixture, framework::DatasetMode::ALL, combine(datasets::SmallShapes(), ElementwisePowerFP32Dataset)) +FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwisePowerFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapes(), ElementwisePowerFP32Dataset), + EmptyActivationFunctionsDataset)) +{ + // Validate output + validate(CLAccessor(_target), _reference, tolerance_fp32); +} +FIXTURE_DATA_TEST_CASE(RunWithActivation, CLElementwisePowerFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::TinyShapes(), ElementwisePowerFP32Dataset), + ActivationFunctionsDataset)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp32); } -FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, CLElementwisePowerBroadcastFixture, framework::DatasetMode::ALL, combine(datasets::SmallShapesBroadcast(), - ElementwisePowerFP32Dataset)) +FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, CLElementwisePowerBroadcastFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapesBroadcast(), + ElementwisePowerFP32Dataset), + EmptyActivationFunctionsDataset)) +{ + // Validate output + validate(CLAccessor(_target), _reference, tolerance_fp32); +} +FIXTURE_DATA_TEST_CASE(RunWithActivationBroadcast, CLElementwisePowerBroadcastFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::TinyShapesBroadcast(), + ElementwisePowerFP32Dataset), + ActivationFunctionsDataset)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp32); diff --git a/tests/validation/CL/ElementwiseSquaredDiff.cpp b/tests/validation/CL/ElementwiseSquaredDiff.cpp index edc150109e..86fdc21d6d 100644 --- a/tests/validation/CL/ElementwiseSquaredDiff.cpp +++ b/tests/validation/CL/ElementwiseSquaredDiff.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2019 ARM Limited. + * Copyright (c) 2018-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -64,6 +64,13 @@ const auto ElementwiseSquaredDiffFP16Dataset = combine(combine(framework::datase framework::dataset::make("DataType", DataType::F16)); const auto ElementwiseSquaredDiffFP32Dataset = combine(combine(framework::dataset::make("DataType", DataType::F32), framework::dataset::make("DataType", DataType::F32)), framework::dataset::make("DataType", DataType::F32)); +const auto EmptyActivationFunctionsDataset = framework::dataset::make("ActivationInfo", +{ ActivationLayerInfo() }); +const auto ActivationFunctionsDataset = framework::dataset::make("ActivationInfo", +{ + ActivationLayerInfo(ActivationLayerInfo::ActivationFunction::BOUNDED_RELU, 0.75f, 0.25f), + ActivationLayerInfo(ActivationLayerInfo::ActivationFunction::LOGISTIC, 0.75f, 0.25f) +}); } // namespace TEST_SUITE(CL) @@ -239,9 +246,19 @@ FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseSquaredDiffQuantizedFixture +using CLElementwiseSquaredDiffFloatFixture = ElementwiseSquaredDiffValidationFloatFixture; + TEST_SUITE(Float) TEST_SUITE(FP16) -FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseSquaredDiffFixture, framework::DatasetMode::ALL, combine(datasets::SmallShapes(), ElementwiseSquaredDiffFP16Dataset)) +FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseSquaredDiffFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapes(), ElementwiseSquaredDiffFP16Dataset), + EmptyActivationFunctionsDataset)) +{ + // Validate output + validate(CLAccessor(_target), _reference, tolerance_fp16, 0.01); +} +FIXTURE_DATA_TEST_CASE(RunWithActivation, CLElementwiseSquaredDiffFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::TinyShapes(), ElementwiseSquaredDiffFP16Dataset), + ActivationFunctionsDataset)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp16, 0.01); @@ -272,16 +289,31 @@ DATA_TEST_CASE(Configuration, framework::DatasetMode::ALL, datasets::SmallShapes validate(dst.info()->padding(), padding); } -FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseSquaredDiffFixture, framework::DatasetMode::ALL, combine(datasets::SmallShapes(), ElementwiseSquaredDiffFP32Dataset)) +FIXTURE_DATA_TEST_CASE(RunSmall, CLElementwiseSquaredDiffFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapes(), ElementwiseSquaredDiffFP32Dataset), + EmptyActivationFunctionsDataset)) +{ + // Validate output + validate(CLAccessor(_target), _reference, tolerance_fp32); +} +FIXTURE_DATA_TEST_CASE(RunWithActivation, CLElementwiseSquaredDiffFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::TinyShapes(), ElementwiseSquaredDiffFP32Dataset), + ActivationFunctionsDataset)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp32); } template -using CLElementwiseSquaredDiffBroadcastFixture = ElementwiseSquaredDiffBroadcastValidationFixture; +using CLElementwiseSquaredDiffBroadcastFloatFixture = ElementwiseSquaredDiffBroadcastValidationFloatFixture; -FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, CLElementwiseSquaredDiffBroadcastFixture, framework::DatasetMode::ALL, combine(datasets::SmallShapesBroadcast(), - ElementwiseSquaredDiffFP32Dataset)) +FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, CLElementwiseSquaredDiffBroadcastFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::SmallShapesBroadcast(), + ElementwiseSquaredDiffFP32Dataset), + EmptyActivationFunctionsDataset)) +{ + // Validate output + validate(CLAccessor(_target), _reference, tolerance_fp32); +} +FIXTURE_DATA_TEST_CASE(RunWithActivationBroadcast, CLElementwiseSquaredDiffBroadcastFloatFixture, framework::DatasetMode::ALL, combine(combine(datasets::TinyShapesBroadcast(), + ElementwiseSquaredDiffFP32Dataset), + ActivationFunctionsDataset)) { // Validate output validate(CLAccessor(_target), _reference, tolerance_fp32); diff --git a/tests/validation/CL/PixelWiseMultiplication.cpp b/tests/validation/CL/PixelWiseMultiplication.cpp index ff9101a997..310828c48d 100644 --- a/tests/validation/CL/PixelWiseMultiplication.cpp +++ b/tests/validation/CL/PixelWiseMultiplication.cpp @@ -43,20 +43,27 @@ namespace const float scale_255 = 1.f / 255.f; constexpr AbsoluteTolerance tolerance_qasymm8(1); /**< Tolerance value for comparing reference's output against implementation's output for 8-bit quantized asymmetric data types */ constexpr AbsoluteTolerance tolerance_qsymm16(1); /**< Tolerance value for comparing reference's output against implementation's output for 16-bit quantized symmetric data types */ +const auto EmptyActivationFunctionsDataset = framework::dataset::make("ActivationInfo", +{ ActivationLayerInfo() }); +const auto ActivationFunctionsDataset = framework::dataset::make("ActivationInfo", +{ + ActivationLayerInfo(ActivationLayerInfo::ActivationFunction::BOUNDED_RELU, 0.75f, 0.25f), + ActivationLayerInfo(ActivationLayerInfo::ActivationFunction::LOGISTIC, 0.75f, 0.25f) +}); } //namespace // *INDENT-OFF* // clang-format off #define VALIDATE(TYPE, TOLERANCE) validate(CLAccessor(_target), _reference, AbsoluteTolerance(TOLERANCE), 0.f); -#define PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(TEST_NAME, FIXTURE, MODE, SHAPES, DT1, DT2, SCALE, RP, VALIDATE) \ +#define PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(TEST_NAME, FIXTURE, MODE, SHAPES, DT1, DT2, SCALE, RP, ACT, VALIDATE) \ FIXTURE_DATA_TEST_CASE(TEST_NAME, CLPixelWiseMultiplication##FIXTURE, framework::DatasetMode::MODE, \ - combine(combine(combine(combine(combine( \ + combine(combine(combine(combine(combine(combine( \ datasets::SHAPES, \ framework::dataset::make("DataType1", DataType::DT1)), \ framework::dataset::make("DataType2", DataType::DT2)), \ framework::dataset::make("Scale", std::move(SCALE))), \ datasets::ConvertPolicies()), \ - framework::dataset::make("RoundingPolicy", RoundingPolicy::RP))) \ + framework::dataset::make("RoundingPolicy", RoundingPolicy::RP)), ACT)) \ { \ VALIDATE \ } @@ -65,11 +72,11 @@ constexpr AbsoluteTolerance tolerance_qsymm16(1); /**< Tolerance value fo } // namespace template -using CLPixelWiseMultiplicationToF16Fixture = PixelWiseMultiplicationValidationFixture; +using CLPixelWiseMultiplicationToF16Fixture = PixelWiseMultiplicationValidationFloatFixture; template -using CLPixelWiseMultiplicationToF32Fixture = PixelWiseMultiplicationValidationFixture; +using CLPixelWiseMultiplicationToF32Fixture = PixelWiseMultiplicationValidationFloatFixture; template -using CLPixelWiseMultiplicationBroadcastFixture = PixelWiseMultiplicationBroadcastValidationFixture; +using CLPixelWiseMultiplicationBroadcastFixture = PixelWiseMultiplicationBroadcastValidationFloatFixture; TEST_SUITE(CL) TEST_SUITE(PixelWiseMultiplication) @@ -110,17 +117,24 @@ DATA_TEST_CASE(Validate, framework::DatasetMode::ALL, zip(zip(zip(zip( TEST_SUITE(F16toF16) TEST_SUITE(Scale255) -PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(RunSmall, ToF16Fixture, PRECOMMIT, SmallShapes(), F16, F16, scale_255, TO_NEAREST_UP, VALIDATE(float, 1.f)) +PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(RunSmall, ToF16Fixture, PRECOMMIT, SmallShapes(), F16, F16, scale_255, TO_NEAREST_UP, EmptyActivationFunctionsDataset, + VALIDATE(float, 1.f)) +PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(RunWithActivation, ToF16Fixture, ALL, TinyShapes(), F16, F16, scale_255, TO_NEAREST_UP, ActivationFunctionsDataset, VALIDATE(float, + 1.f)) TEST_SUITE_END() // Scale255 TEST_SUITE_END() // F16toF16 TEST_SUITE(F32toF32) TEST_SUITE(Scale255) -PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(RunSmall, ToF32Fixture, PRECOMMIT, SmallShapes(), F32, F32, scale_255, TO_NEAREST_UP, VALIDATE(float, 1.f)) +PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(RunSmall, ToF32Fixture, PRECOMMIT, SmallShapes(), F32, F32, scale_255, TO_NEAREST_UP, EmptyActivationFunctionsDataset, VALIDATE(float, 1.f)) +PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(RunWithActivation, ToF32Fixture, ALL, TinyShapes(), F32, F32, scale_255, TO_NEAREST_UP, ActivationFunctionsDataset, VALIDATE(float, 1.f)) TEST_SUITE_END() // Scale255 TEST_SUITE_END() // F32toF32 -PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, BroadcastFixture, PRECOMMIT, SmallShapesBroadcast(), F32, F32, scale_255, TO_NEAREST_UP, VALIDATE(float, 1.f)) +PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, BroadcastFixture, PRECOMMIT, SmallShapesBroadcast(), F32, F32, scale_255, TO_NEAREST_UP, EmptyActivationFunctionsDataset, + VALIDATE(float, 1.f)) +PIXEL_WISE_MULTIPLICATION_FIXTURE_DATA_TEST_CASE(RunWithActivationSmallBroadcast, BroadcastFixture, ALL, TinyShapesBroadcast(), F32, F32, scale_255, TO_NEAREST_UP, ActivationFunctionsDataset, + VALIDATE(float, 1.f)) template using CLPixelWiseMultiplicationQuantizedFixture = PixelWiseMultiplicationValidationQuantizedFixture; diff --git a/tests/validation/NEON/ElementwiseDivision.cpp b/tests/validation/NEON/ElementwiseDivision.cpp index e68ce19eed..f5e1f86dbc 100644 --- a/tests/validation/NEON/ElementwiseDivision.cpp +++ b/tests/validation/NEON/ElementwiseDivision.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019 ARM Limited. + * Copyright (c) 2019-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -57,7 +57,7 @@ TEST_SUITE(NEON) TEST_SUITE(ElementwiseDivision) template -using NEElementwiseDivisionFixture = ElementwiseDivisionValidationFixture; +using NEElementwiseDivisionFixture = ArithmeticDivisionValidationFixture; // *INDENT-OFF* // clang-format off @@ -124,7 +124,7 @@ FIXTURE_DATA_TEST_CASE(RunSmall, NEElementwiseDivisionFixture, framework: } template -using NEElementwiseDivisionBroadcastFixture = ElementwiseDivisionBroadcastValidationFixture; +using NEElementwiseDivisionBroadcastFixture = ArithmeticDivisionBroadcastValidationFixture; FIXTURE_DATA_TEST_CASE(RunSmallBroadcast, NEElementwiseDivisionBroadcastFixture, framework::DatasetMode::ALL, combine(datasets::SmallShapesBroadcast(), ElementwiseDivisionFP32Dataset)) diff --git a/tests/validation/fixtures/ArithmeticOperationsFixture.h b/tests/validation/fixtures/ArithmeticOperationsFixture.h index d495ab1049..4a6b0bd3f3 100644 --- a/tests/validation/fixtures/ArithmeticOperationsFixture.h +++ b/tests/validation/fixtures/ArithmeticOperationsFixture.h @@ -32,6 +32,7 @@ #include "tests/framework/Asserts.h" #include "tests/framework/Fixture.h" #include "tests/validation/Helpers.h" +#include "tests/validation/reference/ActivationLayer.h" #include "tests/validation/reference/ArithmeticOperations.h" namespace arm_compute @@ -47,9 +48,10 @@ public: template void setup(reference::ArithmeticOperation op, const TensorShape &shape0, const TensorShape &shape1, DataType data_type0, DataType data_type1, DataType output_data_type, ConvertPolicy convert_policy, - QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out) + QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out, ActivationLayerInfo act_info) { _op = op; + _act_info = act_info; _target = compute_target(shape0, shape1, data_type0, data_type1, output_data_type, convert_policy, qinfo0, qinfo1, qinfo_out); _reference = compute_reference(shape0, shape1, data_type0, data_type1, output_data_type, convert_policy, qinfo0, qinfo1, qinfo_out); } @@ -71,7 +73,7 @@ protected: // Create and configure function FunctionType arith_op; - arith_op.configure(&ref_src1, &ref_src2, &dst, convert_policy); + arith_op.configure(&ref_src1, &ref_src2, &dst, convert_policy, _act_info); ARM_COMPUTE_EXPECT(ref_src1.info()->is_resizable(), framework::LogLevel::ERRORS); ARM_COMPUTE_EXPECT(ref_src2.info()->is_resizable(), framework::LogLevel::ERRORS); @@ -109,12 +111,14 @@ protected: fill(ref_src1, 0); fill(ref_src2, 1); - return reference::arithmetic_operation(_op, ref_src1, ref_src2, ref_dst, convert_policy); + auto result = reference::arithmetic_operation(_op, ref_src1, ref_src2, ref_dst, convert_policy); + return _act_info.enabled() ? reference::activation_layer(result, _act_info, qinfo_out) : result; } TensorType _target{}; SimpleTensor _reference{}; reference::ArithmeticOperation _op{ reference::ArithmeticOperation::ADD }; + ActivationLayerInfo _act_info{}; }; template @@ -125,7 +129,7 @@ public: void setup(const TensorShape &shape0, const TensorShape &shape1, DataType data_type0, DataType data_type1, DataType output_data_type, ConvertPolicy convert_policy) { ArithmeticOperationGenericFixture::setup(reference::ArithmeticOperation::ADD, shape0, shape1, data_type0, data_type1, - output_data_type, convert_policy, QuantizationInfo(), QuantizationInfo(), QuantizationInfo()); + output_data_type, convert_policy, QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), ActivationLayerInfo()); } }; @@ -137,7 +141,31 @@ public: void setup(const TensorShape &shape, DataType data_type0, DataType data_type1, DataType output_data_type, ConvertPolicy convert_policy) { ArithmeticOperationGenericFixture::setup(reference::ArithmeticOperation::ADD, shape, shape, data_type0, data_type1, - output_data_type, convert_policy, QuantizationInfo(), QuantizationInfo(), QuantizationInfo()); + output_data_type, convert_policy, QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), ActivationLayerInfo()); + } +}; + +template +class ArithmeticAdditionBroadcastValidationFloatFixture : public ArithmeticOperationGenericFixture +{ +public: + template + void setup(const TensorShape &shape0, const TensorShape &shape1, DataType data_type0, DataType data_type1, DataType output_data_type, ConvertPolicy convert_policy, ActivationLayerInfo act_info) + { + ArithmeticOperationGenericFixture::setup(reference::ArithmeticOperation::ADD, shape0, shape1, data_type0, data_type1, + output_data_type, convert_policy, QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), act_info); + } +}; + +template +class ArithmeticAdditionValidationFloatFixture : public ArithmeticOperationGenericFixture +{ +public: + template + void setup(const TensorShape &shape, DataType data_type0, DataType data_type1, DataType output_data_type, ConvertPolicy convert_policy, ActivationLayerInfo act_info) + { + ArithmeticOperationGenericFixture::setup(reference::ArithmeticOperation::ADD, shape, shape, data_type0, data_type1, + output_data_type, convert_policy, QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), act_info); } }; @@ -151,7 +179,7 @@ public: { ArithmeticOperationGenericFixture::setup(reference::ArithmeticOperation::ADD, shape, shape, data_type0, data_type1, - output_data_type, convert_policy, qinfo0, qinfo1, qinfo_out); + output_data_type, convert_policy, qinfo0, qinfo1, qinfo_out, ActivationLayerInfo()); } }; @@ -164,7 +192,20 @@ public: { ArithmeticOperationGenericFixture::setup(reference::ArithmeticOperation::SUB, shape0, shape1, data_type0, data_type1, output_data_type, convert_policy, - QuantizationInfo(), QuantizationInfo(), QuantizationInfo()); + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), ActivationLayerInfo()); + } +}; + +template +class ArithmeticSubtractionBroadcastValidationFloatFixture : public ArithmeticOperationGenericFixture +{ +public: + template + void setup(const TensorShape &shape0, const TensorShape &shape1, DataType data_type0, DataType data_type1, DataType output_data_type, ConvertPolicy convert_policy, ActivationLayerInfo act_info) + { + ArithmeticOperationGenericFixture::setup(reference::ArithmeticOperation::SUB, shape0, shape1, + data_type0, data_type1, output_data_type, convert_policy, + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), act_info); } }; @@ -177,7 +218,20 @@ public: { ArithmeticOperationGenericFixture::setup(reference::ArithmeticOperation::SUB, shape, shape, data_type0, data_type1, output_data_type, convert_policy, - QuantizationInfo(), QuantizationInfo(), QuantizationInfo()); + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), ActivationLayerInfo()); + } +}; + +template +class ArithmeticSubtractionValidationFloatFixture : public ArithmeticOperationGenericFixture +{ +public: + template + void setup(const TensorShape &shape, DataType data_type0, DataType data_type1, DataType output_data_type, ConvertPolicy convert_policy, ActivationLayerInfo act_info) + { + ArithmeticOperationGenericFixture::setup(reference::ArithmeticOperation::SUB, shape, shape, + data_type0, data_type1, output_data_type, convert_policy, + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), act_info); } }; @@ -192,7 +246,7 @@ public: { ArithmeticOperationGenericFixture::setup(reference::ArithmeticOperation::SUB, shape, shape, data_type0, data_type1, output_data_type, - convert_policy, qinfo0, qinfo1, qinfo_out); + convert_policy, qinfo0, qinfo1, qinfo_out, ActivationLayerInfo()); } }; } // namespace validation diff --git a/tests/validation/fixtures/ElementwiseOperationsFixture.h b/tests/validation/fixtures/ElementwiseOperationsFixture.h index de61c487e6..44c096c521 100644 --- a/tests/validation/fixtures/ElementwiseOperationsFixture.h +++ b/tests/validation/fixtures/ElementwiseOperationsFixture.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2018-2019 ARM Limited. + * Copyright (c) 2018-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -32,6 +32,7 @@ #include "tests/framework/Asserts.h" #include "tests/framework/Fixture.h" #include "tests/validation/Helpers.h" +#include "tests/validation/reference/ActivationLayer.h" #include "tests/validation/reference/ElementwiseOperations.h" namespace arm_compute @@ -127,6 +128,70 @@ protected: ArithmeticOperation _op{ ArithmeticOperation::ADD }; }; +// Arithmetic operation fused with activation function +template +class ArithmeticOperationsFuseActivationFixture : public ArithmeticOperationsGenericFixture +{ +public: + template + void setup(ArithmeticOperation op, const TensorShape &shape0, const TensorShape &shape1, + DataType data_type0, DataType data_type1, DataType output_data_type, + QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out, ActivationLayerInfo act_info) + { + ArithmeticOperationsGenericFixture::setup(op, shape0, shape1, + data_type0, data_type1, output_data_type, + qinfo0, qinfo1, qinfo_out); + _act_info = act_info; + } + +protected: + TensorType compute_target(const TensorShape &shape0, const TensorShape &shape1, DataType data_type0, DataType data_type1, DataType output_data_type, + QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out) + { + // Create tensors + TensorType ref_src1 = create_tensor(shape0, data_type0, 1, qinfo0); + TensorType ref_src2 = create_tensor(shape1, data_type1, 1, qinfo1); + TensorType dst = create_tensor(TensorShape::broadcast_shape(shape0, shape1), output_data_type, 1, qinfo_out); + + // Create and configure function + FunctionType elem_op; + elem_op.configure(&ref_src1, &ref_src2, &dst, _act_info); + + ARM_COMPUTE_EXPECT(ref_src1.info()->is_resizable(), framework::LogLevel::ERRORS); + ARM_COMPUTE_EXPECT(ref_src2.info()->is_resizable(), framework::LogLevel::ERRORS); + ARM_COMPUTE_EXPECT(dst.info()->is_resizable(), framework::LogLevel::ERRORS); + + // Allocate tensors + ref_src1.allocator()->allocate(); + ref_src2.allocator()->allocate(); + dst.allocator()->allocate(); + + ARM_COMPUTE_EXPECT(!ref_src1.info()->is_resizable(), framework::LogLevel::ERRORS); + ARM_COMPUTE_EXPECT(!ref_src2.info()->is_resizable(), framework::LogLevel::ERRORS); + ARM_COMPUTE_EXPECT(!dst.info()->is_resizable(), framework::LogLevel::ERRORS); + + // Fill tensors + fill(AccessorType(ref_src1), 0); + fill(AccessorType(ref_src2), 1); + + // Compute function + elem_op.run(); + + return dst; + } + + SimpleTensor compute_reference(const TensorShape &shape0, const TensorShape &shape1, + DataType data_type0, DataType data_type1, DataType output_data_type, + QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out) + { + auto result = ArithmeticOperationsGenericFixture::compute_reference(shape0, shape1, data_type0, + data_type1, output_data_type, qinfo0, qinfo1, qinfo_out); + return _act_info.enabled() ? reference::activation_layer(result, _act_info, qinfo_out) : result; + } + + ActivationLayerInfo _act_info{}; +}; + template class ArithmeticDivisionBroadcastValidationFixture : public ArithmeticOperationsGenericFixture { @@ -153,6 +218,32 @@ public: } }; +template +class ArithmeticDivisionBroadcastValidationFloatFixture : public ArithmeticOperationsFuseActivationFixture +{ +public: + template + void setup(const TensorShape &shape0, const TensorShape &shape1, DataType data_type0, DataType data_type1, DataType output_data_type, ActivationLayerInfo act_info) + { + ArithmeticOperationsFuseActivationFixture::setup(ArithmeticOperation::DIV, shape0, shape1, + data_type0, data_type1, output_data_type, + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), act_info); + } +}; + +template +class ArithmeticDivisionValidationFloatFixture : public ArithmeticOperationsFuseActivationFixture +{ +public: + template + void setup(const TensorShape &shape, DataType data_type0, DataType data_type1, DataType output_data_type, ActivationLayerInfo act_info) + { + ArithmeticOperationsFuseActivationFixture::setup(ArithmeticOperation::DIV, shape, shape, + data_type0, data_type1, output_data_type, + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), act_info); + } +}; + template class ArithmeticDivisionValidationQuantizedFixture : public ArithmeticOperationsGenericFixture { @@ -194,6 +285,32 @@ public: } }; +template +class ElementwiseMaxBroadcastValidationFloatFixture : public ArithmeticOperationsFuseActivationFixture +{ +public: + template + void setup(const TensorShape &shape0, const TensorShape &shape1, DataType data_type0, DataType data_type1, DataType output_data_type, ActivationLayerInfo act_info) + { + ArithmeticOperationsFuseActivationFixture::setup(ArithmeticOperation::MAX, shape0, shape1, + data_type0, data_type1, output_data_type, + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), act_info); + } +}; + +template +class ElementwiseMaxValidationFloatFixture : public ArithmeticOperationsFuseActivationFixture +{ +public: + template + void setup(const TensorShape &shape, DataType data_type0, DataType data_type1, DataType output_data_type, ActivationLayerInfo act_info) + { + ArithmeticOperationsFuseActivationFixture::setup(ArithmeticOperation::MAX, shape, shape, + data_type0, data_type1, output_data_type, + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), act_info); + } +}; + template class ElementwiseMaxValidationQuantizedFixture : public ArithmeticOperationsGenericFixture { @@ -250,6 +367,32 @@ public: } }; +template +class ElementwiseMinBroadcastValidationFloatFixture : public ArithmeticOperationsFuseActivationFixture +{ +public: + template + void setup(const TensorShape &shape0, const TensorShape &shape1, DataType data_type0, DataType data_type1, DataType output_data_type, ActivationLayerInfo act_info) + { + ArithmeticOperationsFuseActivationFixture::setup(ArithmeticOperation::MIN, shape0, shape1, + data_type0, data_type1, output_data_type, + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), act_info); + } +}; + +template +class ElementwiseMinValidationFloatFixture : public ArithmeticOperationsFuseActivationFixture +{ +public: + template + void setup(const TensorShape &shape, DataType data_type0, DataType data_type1, DataType output_data_type, ActivationLayerInfo act_info) + { + ArithmeticOperationsFuseActivationFixture::setup(ArithmeticOperation::MIN, shape, shape, + data_type0, data_type1, output_data_type, + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), act_info); + } +}; + template class ElementwiseMinValidationQuantizedFixture : public ArithmeticOperationsGenericFixture { @@ -306,6 +449,32 @@ public: } }; +template +class ElementwiseSquaredDiffBroadcastValidationFloatFixture : public ArithmeticOperationsFuseActivationFixture +{ +public: + template + void setup(const TensorShape &shape0, const TensorShape &shape1, DataType data_type0, DataType data_type1, DataType output_data_type, ActivationLayerInfo act_info) + { + ArithmeticOperationsFuseActivationFixture::setup(ArithmeticOperation::SQUARED_DIFF, shape0, shape1, + data_type0, data_type1, output_data_type, + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), act_info); + } +}; + +template +class ElementwiseSquaredDiffValidationFloatFixture : public ArithmeticOperationsFuseActivationFixture +{ +public: + template + void setup(const TensorShape &shape, DataType data_type0, DataType data_type1, DataType output_data_type, ActivationLayerInfo act_info) + { + ArithmeticOperationsFuseActivationFixture::setup(ArithmeticOperation::SQUARED_DIFF, shape, shape, + data_type0, data_type1, output_data_type, + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), act_info); + } +}; + template class ElementwiseSquaredDiffValidationQuantizedFixture : public ArithmeticOperationsGenericFixture { @@ -393,84 +562,54 @@ public: }; template -class ElementwiseDivisionBroadcastValidationFixture : public ArithmeticOperationsGenericFixture +class ElementwisePowerBroadcastValidationFixture : public ArithmeticOperationsGenericFixture { public: template void setup(const TensorShape &shape0, const TensorShape &shape1, DataType data_type0, DataType data_type1, DataType output_data_type) { - ArithmeticOperationsGenericFixture::setup(ArithmeticOperation::DIV, shape0, shape1, + ArithmeticOperationsGenericFixture::setup(ArithmeticOperation::POWER, shape0, shape1, data_type0, data_type1, output_data_type, QuantizationInfo(), QuantizationInfo(), QuantizationInfo()); } }; template -class ElementwiseDivisionValidationFixture : public ArithmeticOperationsGenericFixture +class ElementwisePowerValidationFixture : public ArithmeticOperationsGenericFixture { public: template void setup(const TensorShape &shape, DataType data_type0, DataType data_type1, DataType output_data_type) { - ArithmeticOperationsGenericFixture::setup(ArithmeticOperation::DIV, shape, shape, + ArithmeticOperationsGenericFixture::setup(ArithmeticOperation::POWER, shape, shape, data_type0, data_type1, output_data_type, QuantizationInfo(), QuantizationInfo(), QuantizationInfo()); } }; template -class ElementwiseDivisionValidationQuantizedFixture : public ArithmeticOperationsGenericFixture -{ -public: - template - void setup(const TensorShape &shape, DataType data_type0, DataType data_type1, DataType output_data_type, - QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out) - - { - ArithmeticOperationsGenericFixture::setup(ArithmeticOperation::DIV, shape, shape, - data_type0, data_type1, output_data_type, - qinfo0, qinfo1, qinfo_out); - } -}; - -template -class ElementwiseDivisionQuantizedBroadcastValidationFixture : public ArithmeticOperationsGenericFixture -{ -public: - template - void setup(const TensorShape &shape0, const TensorShape &shape1, DataType data_type0, DataType data_type1, DataType output_data_type, - QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out) - - { - ArithmeticOperationsGenericFixture::setup(ArithmeticOperation::DIV, shape0, shape1, - data_type0, data_type1, output_data_type, - qinfo0, qinfo1, qinfo_out); - } -}; - -template -class ElementwisePowerBroadcastValidationFixture : public ArithmeticOperationsGenericFixture +class ElementwisePowerBroadcastValidationFloatFixture : public ArithmeticOperationsFuseActivationFixture { public: template - void setup(const TensorShape &shape0, const TensorShape &shape1, DataType data_type0, DataType data_type1, DataType output_data_type) + void setup(const TensorShape &shape0, const TensorShape &shape1, DataType data_type0, DataType data_type1, DataType output_data_type, ActivationLayerInfo act_info) { - ArithmeticOperationsGenericFixture::setup(ArithmeticOperation::POWER, shape0, shape1, - data_type0, data_type1, output_data_type, - QuantizationInfo(), QuantizationInfo(), QuantizationInfo()); + ArithmeticOperationsFuseActivationFixture::setup(ArithmeticOperation::POWER, shape0, shape1, + data_type0, data_type1, output_data_type, + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), act_info); } }; template -class ElementwisePowerValidationFixture : public ArithmeticOperationsGenericFixture +class ElementwisePowerValidationFloatFixture : public ArithmeticOperationsFuseActivationFixture { public: template - void setup(const TensorShape &shape, DataType data_type0, DataType data_type1, DataType output_data_type) + void setup(const TensorShape &shape, DataType data_type0, DataType data_type1, DataType output_data_type, ActivationLayerInfo act_info) { - ArithmeticOperationsGenericFixture::setup(ArithmeticOperation::POWER, shape, shape, - data_type0, data_type1, output_data_type, - QuantizationInfo(), QuantizationInfo(), QuantizationInfo()); + ArithmeticOperationsFuseActivationFixture::setup(ArithmeticOperation::POWER, shape, shape, + data_type0, data_type1, output_data_type, + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), act_info); } }; diff --git a/tests/validation/fixtures/PixelWiseMultiplicationFixture.h b/tests/validation/fixtures/PixelWiseMultiplicationFixture.h index 37359f421b..f561a37a71 100644 --- a/tests/validation/fixtures/PixelWiseMultiplicationFixture.h +++ b/tests/validation/fixtures/PixelWiseMultiplicationFixture.h @@ -31,6 +31,7 @@ #include "tests/IAccessor.h" #include "tests/framework/Asserts.h" #include "tests/framework/Fixture.h" +#include "tests/validation/reference/ActivationLayer.h" #include "tests/validation/reference/PixelWiseMultiplication.h" namespace arm_compute @@ -46,18 +47,19 @@ public: template void setup(const TensorShape &shape0, const TensorShape &shape1, - DataType dt_in1, - DataType dt_in2, - DataType dt_out, - float scale, - ConvertPolicy convert_policy, - RoundingPolicy rounding_policy, - QuantizationInfo qinfo0, - QuantizationInfo qinfo1, - QuantizationInfo qinfo_out) + DataType dt_in1, + DataType dt_in2, + DataType dt_out, + float scale, + ConvertPolicy convert_policy, + RoundingPolicy rounding_policy, + QuantizationInfo qinfo0, + QuantizationInfo qinfo1, + QuantizationInfo qinfo_out, + ActivationLayerInfo act_info) { - _target = compute_target(shape0, shape1, dt_in1, dt_in2, dt_out, scale, convert_policy, rounding_policy, qinfo0, qinfo1, qinfo_out); - _reference = compute_reference(shape0, shape1, dt_in1, dt_in2, dt_out, scale, convert_policy, rounding_policy, qinfo0, qinfo1, qinfo_out); + _target = compute_target(shape0, shape1, dt_in1, dt_in2, dt_out, scale, convert_policy, rounding_policy, qinfo0, qinfo1, qinfo_out, act_info); + _reference = compute_reference(shape0, shape1, dt_in1, dt_in2, dt_out, scale, convert_policy, rounding_policy, qinfo0, qinfo1, qinfo_out, act_info); } protected: @@ -69,7 +71,7 @@ protected: TensorType compute_target(const TensorShape &shape0, const TensorShape &shape1, DataType dt_in1, DataType dt_in2, DataType dt_out, float scale, ConvertPolicy convert_policy, RoundingPolicy rounding_policy, - QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out) + QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out, ActivationLayerInfo act_info) { // Create tensors TensorType src1 = create_tensor(shape0, dt_in1, 1, qinfo0); @@ -78,7 +80,7 @@ protected: // Create and configure function FunctionType multiply; - multiply.configure(&src1, &src2, &dst, scale, convert_policy, rounding_policy); + multiply.configure(&src1, &src2, &dst, scale, convert_policy, rounding_policy, act_info); ARM_COMPUTE_EXPECT(src1.info()->is_resizable(), framework::LogLevel::ERRORS); ARM_COMPUTE_EXPECT(src2.info()->is_resizable(), framework::LogLevel::ERRORS); @@ -105,7 +107,7 @@ protected: SimpleTensor compute_reference(const TensorShape &shape0, const TensorShape &shape1, DataType dt_in1, DataType dt_in2, DataType dt_out, float scale, ConvertPolicy convert_policy, RoundingPolicy rounding_policy, - QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out) + QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out, ActivationLayerInfo act_info) { // Create reference SimpleTensor src1{ shape0, dt_in1, 1, qinfo0 }; @@ -115,7 +117,8 @@ protected: fill(src1, 0); fill(src2, 1); - return reference::pixel_wise_multiplication(src1, src2, scale, convert_policy, rounding_policy, dt_out, qinfo_out); + auto result = reference::pixel_wise_multiplication(src1, src2, scale, convert_policy, rounding_policy, dt_out, qinfo_out); + return act_info.enabled() ? reference::activation_layer(result, act_info, qinfo_out) : result; } TensorType _target{}; @@ -130,7 +133,7 @@ public: void setup(const TensorShape &shape, DataType dt_in1, DataType dt_in2, float scale, ConvertPolicy convert_policy, RoundingPolicy rounding_policy) { PixelWiseMultiplicationGenericValidationFixture::setup(shape, shape, dt_in1, dt_in2, dt_in2, scale, convert_policy, rounding_policy, - QuantizationInfo(), QuantizationInfo(), QuantizationInfo()); + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), ActivationLayerInfo()); } }; @@ -142,7 +145,32 @@ public: void setup(const TensorShape &shape0, const TensorShape &shape1, DataType dt_in1, DataType dt_in2, float scale, ConvertPolicy convert_policy, RoundingPolicy rounding_policy) { PixelWiseMultiplicationGenericValidationFixture::setup(shape0, shape1, dt_in1, dt_in2, dt_in2, scale, convert_policy, rounding_policy, - QuantizationInfo(), QuantizationInfo(), QuantizationInfo()); + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), ActivationLayerInfo()); + } +}; + +template +class PixelWiseMultiplicationValidationFloatFixture : public PixelWiseMultiplicationGenericValidationFixture +{ +public: + template + void setup(const TensorShape &shape, DataType dt_in1, DataType dt_in2, float scale, ConvertPolicy convert_policy, RoundingPolicy rounding_policy, ActivationLayerInfo act_info) + { + PixelWiseMultiplicationGenericValidationFixture::setup(shape, shape, dt_in1, dt_in2, dt_in2, scale, convert_policy, rounding_policy, + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), act_info); + } +}; + +template +class PixelWiseMultiplicationBroadcastValidationFloatFixture : public PixelWiseMultiplicationGenericValidationFixture +{ +public: + template + void setup(const TensorShape &shape0, const TensorShape &shape1, DataType dt_in1, DataType dt_in2, float scale, ConvertPolicy convert_policy, RoundingPolicy rounding_policy, + ActivationLayerInfo act_info) + { + PixelWiseMultiplicationGenericValidationFixture::setup(shape0, shape1, dt_in1, dt_in2, dt_in2, scale, convert_policy, rounding_policy, + QuantizationInfo(), QuantizationInfo(), QuantizationInfo(), act_info); } }; @@ -154,8 +182,8 @@ public: void setup(const TensorShape &shape, DataType dt_in1, DataType dt_in2, DataType dt_out, float scale, ConvertPolicy convert_policy, RoundingPolicy rounding_policy, QuantizationInfo qinfo0, QuantizationInfo qinfo1, QuantizationInfo qinfo_out) { - PixelWiseMultiplicationGenericValidationFixture::setup(shape, shape, dt_in1, dt_in2, dt_out, scale, convert_policy, - rounding_policy, qinfo0, qinfo1, qinfo_out); + PixelWiseMultiplicationGenericValidationFixture::setup(shape, shape, dt_in1, dt_in2, dt_out, scale, convert_policy, rounding_policy, + qinfo0, qinfo1, qinfo_out, ActivationLayerInfo()); } }; } // namespace validation diff --git a/tests/validation/reference/ActivationLayer.cpp b/tests/validation/reference/ActivationLayer.cpp index 7a699c5f86..4aa0f880da 100644 --- a/tests/validation/reference/ActivationLayer.cpp +++ b/tests/validation/reference/ActivationLayer.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017-2019 ARM Limited. + * Copyright (c) 2017-2020 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -87,6 +87,7 @@ SimpleTensor activation_layer(const SimpleTensor &src return dst; } +template SimpleTensor activation_layer(const SimpleTensor &src, ActivationLayerInfo info, const QuantizationInfo &oq_info); template SimpleTensor activation_layer(const SimpleTensor &src, ActivationLayerInfo info, const QuantizationInfo &oq_info); template SimpleTensor activation_layer(const SimpleTensor &src, ActivationLayerInfo info, const QuantizationInfo &oq_info); } // namespace reference diff --git a/tests/validation/reference/ActivationLayer.h b/tests/validation/reference/ActivationLayer.h index f41e87123e..4585a9db10 100644 --- a/tests/validation/reference/ActivationLayer.h +++ b/tests/validation/reference/ActivationLayer.h @@ -82,7 +82,7 @@ inline T activate_float(T x, T a, T b, ActivationLayerInfo::ActivationFunction a ret = x; break; case ActivationLayerInfo::ActivationFunction::HARD_SWISH: - ret = x * ((std::min(std::max((x + 3), 0.0f), 6.0f)) * 0.166666667f); + ret = x * ((std::min(std::max(static_cast(x + 3), static_cast(0.0f)), static_cast(6.0f))) * 0.166666667f); break; default: ARM_COMPUTE_ERROR("Unsupported activation function"); -- cgit v1.2.1