COMPMID-1740: Fuse batch normalization with Convolution Layer at graph level
Change-Id: I77ca51c2c72783cc26a099a6a9c3210cdbbe822d
Signed-off-by: giuros01 <giuseppe.rossini@arm.com>
Reviewed-on: https://review.mlplatform.org/c/797
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Michele Di Giorgio <michele.digiorgio@arm.com>
Reviewed-by: Georgios Pinitas <georgios.pinitas@arm.com>
diff --git a/arm_compute/graph/INodeVisitor.h b/arm_compute/graph/INodeVisitor.h
index 573d642..842ca4b 100644
--- a/arm_compute/graph/INodeVisitor.h
+++ b/arm_compute/graph/INodeVisitor.h
@@ -1,5 +1,5 @@
/*
- * Copyright (c) 2018 ARM Limited.
+ * Copyright (c) 2018-2019 ARM Limited.
*
* SPDX-License-Identifier: MIT
*
@@ -91,6 +91,11 @@
* @param[in] n Node to visit.
*/
virtual void visit(FullyConnectedLayerNode &n) = 0;
+ /** Visit FusedConvolutionBatchNormalizationNode.
+ *
+ * @param[in] n Node to visit.
+ */
+ virtual void visit(FusedConvolutionBatchNormalizationNode &n) = 0;
/** Visit InputNode.
*
* @param[in] n Node to visit.
@@ -195,6 +200,10 @@
{
default_visit();
}
+ virtual void visit(FusedConvolutionBatchNormalizationNode &n) override
+ {
+ default_visit();
+ }
virtual void visit(InputNode &n) override
{
default_visit();
diff --git a/arm_compute/graph/TypePrinter.h b/arm_compute/graph/TypePrinter.h
index ca62d4e..b1cfbcf 100644
--- a/arm_compute/graph/TypePrinter.h
+++ b/arm_compute/graph/TypePrinter.h
@@ -98,6 +98,9 @@
case NodeType::FullyConnectedLayer:
os << "FullyConnectedLayer";
break;
+ case NodeType::FusedConvolutionBatchNormalizationLayer:
+ os << "FusedConvolutionBatchNormalizationLayer";
+ break;
case NodeType::GenerateProposalsLayer:
os << "GenerateProposalsLayer";
break;
diff --git a/arm_compute/graph/Types.h b/arm_compute/graph/Types.h
index 8377253..2905dfcb 100644
--- a/arm_compute/graph/Types.h
+++ b/arm_compute/graph/Types.h
@@ -138,6 +138,7 @@
EltwiseLayer,
FlattenLayer,
FullyConnectedLayer,
+ FusedConvolutionBatchNormalizationLayer,
GenerateProposalsLayer,
NormalizationLayer,
NormalizePlanarYUVLayer,
diff --git a/arm_compute/graph/backends/FunctionHelpers.h b/arm_compute/graph/backends/FunctionHelpers.h
index 7242bc6..d0035d9 100644
--- a/arm_compute/graph/backends/FunctionHelpers.h
+++ b/arm_compute/graph/backends/FunctionHelpers.h
@@ -28,6 +28,7 @@
#include "arm_compute/graph/Tensor.h"
#include "arm_compute/graph/TypePrinter.h"
#include "arm_compute/graph/Types.h"
+#include "arm_compute/graph/backends/FusedConvolutionBatchNormalizationFunction.h"
#include "arm_compute/graph/backends/Utils.h"
#include "arm_compute/graph/nodes/Nodes.h"
@@ -135,11 +136,12 @@
validate_node<TargetInfo>(node, 5 /* expected inputs */, 1 /* expected outputs */);
// Extract IO and info
- typename TargetInfo::TensorType *input = get_backing_tensor<TargetInfo>(node.input(0));
- typename TargetInfo::TensorType *mean = get_backing_tensor<TargetInfo>(node.input(1));
- typename TargetInfo::TensorType *var = get_backing_tensor<TargetInfo>(node.input(2));
- typename TargetInfo::TensorType *beta = get_backing_tensor<TargetInfo>(node.input(3));
- typename TargetInfo::TensorType *gamma = get_backing_tensor<TargetInfo>(node.input(4));
+ typename TargetInfo::TensorType *input = get_backing_tensor<TargetInfo>(node.input(0));
+ typename TargetInfo::TensorType *mean = get_backing_tensor<TargetInfo>(node.input(1));
+ typename TargetInfo::TensorType *var = get_backing_tensor<TargetInfo>(node.input(2));
+ typename TargetInfo::TensorType *beta = get_backing_tensor<TargetInfo>(node.input(3));
+ typename TargetInfo::TensorType *gamma = get_backing_tensor<TargetInfo>(node.input(4));
+
typename TargetInfo::TensorType *output = get_backing_tensor<TargetInfo>(node.output(0));
const float epsilon = node.epsilon();
const ActivationLayerInfo fused_act = node.fused_activation();
@@ -163,6 +165,61 @@
return std::move(func);
}
+/** Create a backend batch normalization layer function
+ *
+ * @tparam BatchNormalizationLayerFunction Backend batch normalization function
+ * @tparam TargetInfo Target-specific information
+ *
+ * @param[in] node Node to create the backend function for
+ *
+ * @return Backend batch normalization layer function
+ */
+template <typename FusedLayerTypes, typename TargetInfo>
+std::unique_ptr<IFunction> create_fused_convolution_batch_normalization_layer(FusedConvolutionBatchNormalizationNode &node)
+{
+ validate_node<TargetInfo>(node, 7 /* expected inputs */, 1 /* expected outputs */);
+
+ // Extract IO and info
+ typename TargetInfo::TensorType *input = get_backing_tensor<TargetInfo>(node.input(0));
+ typename TargetInfo::TensorType *weights = get_backing_tensor<TargetInfo>(node.input(1));
+ typename TargetInfo::TensorType *biases = get_backing_tensor<TargetInfo>(node.input(2));
+ typename TargetInfo::TensorType *mean = get_backing_tensor<TargetInfo>(node.input(3));
+ typename TargetInfo::TensorType *var = get_backing_tensor<TargetInfo>(node.input(4));
+ typename TargetInfo::TensorType *beta = get_backing_tensor<TargetInfo>(node.input(5));
+ typename TargetInfo::TensorType *gamma = get_backing_tensor<TargetInfo>(node.input(6));
+
+ typename TargetInfo::TensorType *output = get_backing_tensor<TargetInfo>(node.output(0));
+
+ const PadStrideInfo conv_info = node.convolution_info();
+ const unsigned int num_groups = node.num_groups();
+ const bool fast_math = node.fast_math_hint() == FastMathHint::Enabled;
+ const ActivationLayerInfo fused_act = node.fused_activation();
+ const float epsilon = node.epsilon();
+
+ const bool is_quantized = is_data_type_quantized_asymmetric(input->info()->data_type());
+ if(is_quantized && biases != nullptr)
+ {
+ biases->info()->set_data_type(DataType::S32);
+ }
+
+ // Create and configure function
+ auto func = support::cpp14::make_unique<FusedConvolutionBatchNormalizationFunction<TargetInfo, FusedLayerTypes>>();
+ func->configure(input, weights, biases, output, mean, var, beta, gamma, epsilon, conv_info, num_groups, fast_math, fused_act);
+
+ // Log info
+ ARM_COMPUTE_LOG_GRAPH_INFO("Instantiated "
+ << node.name()
+ << " Type: " << node.name()
+ << " Target: " << TargetInfo::TargetType
+ << " Data Type: " << input->info()->data_type()
+ << " Input shape: " << input->info()->tensor_shape()
+ << " Weights shape: " << weights->info()->tensor_shape()
+ << " Output shape: " << output->info()->tensor_shape()
+ << (fused_act.enabled() ? " " + to_string(fused_act.activation()) : "")
+ << std::endl);
+ return std::move(func);
+}
+
/** Create a backend bounding box transform layer function
*
* @tparam BoundingBoxTransformLayerFunction Backend bounding box transform function
diff --git a/arm_compute/graph/backends/FusedConvolutionBatchNormalizationFunction.h b/arm_compute/graph/backends/FusedConvolutionBatchNormalizationFunction.h
new file mode 100644
index 0000000..92af17b
--- /dev/null
+++ b/arm_compute/graph/backends/FusedConvolutionBatchNormalizationFunction.h
@@ -0,0 +1,133 @@
+/*
+ * Copyright (c) 2019 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+
+#ifndef __ARM_COMPUTE_GRAPH_BACKENDS_FUSED_CONVOLUTION_BATCH_NORMAZLIZATION_FUNCTION_H__
+#define __ARM_COMPUTE_GRAPH_BACKENDS_FUSED_CONVOLUTION_BATCH_NORMAZLIZATION_FUNCTION_H__
+
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/IFunction.h"
+
+namespace arm_compute
+{
+namespace graph
+{
+namespace backends
+{
+/** Wrapper function to first apply {NE, CL}BatchNormalizationLayer on the weights and then run {NE, CL}ConvolutionLayer with the modified weights */
+template <typename TargetInfo, typename FusedLayerTypes>
+class FusedConvolutionBatchNormalizationFunction : public IFunction
+{
+public:
+ using TensorType = typename TargetInfo::TensorType;
+ using TensorConcreteType = typename TargetInfo::TensorConcreteType;
+
+ FusedConvolutionBatchNormalizationFunction()
+ : _conv_layer(), _fused_batch_norm_layer(), _fused_bias(), _is_prepared(false)
+ {
+ }
+
+ /** Set the input and output tensors.
+ *
+ * @param[in] input Source tensor. 3 lower dimensions represent a single input [width, height, IFM],
+ * while every optional dimension from 4 and above represent a batch of inputs.
+ * Data types supported: QASYMM8/F16/F32.
+ * @param[in] weights Weights tensor. Weights are 4D tensor with dimensions [kernel_x, kernel_y, IFM, OFM]. Data type supported: Same as @p input.
+ * @param[in] bias Biases tensor. Shared biases supported. Biases are 1D tensor with dimensions [OFM].
+ * Data type supported: Should match @p input data type, except for input of QASYMM8 type where biases should be of S32 type.
+ * @param[out] output Destination tensor. 3 lower dimensions represent a single output [width, height, OFM], while the rest represent batch of outputs.
+ * Data types supported: Same as @p input.
+ * @param[in] mean Mean values tensor. 1 dimension with size equal to the feature maps [FM]. Data types supported: Same as @p input
+ * @param[in] var Variance values tensor. 1 dimension with size equal to the feature maps [FM]. Data types supported: Same as @p input
+ * @param[in] beta Beta values tensor info. 1 dimension with size equal to the feature maps [FM]. If not provided, default value for beta is 0. Data types supported: Same as @p input
+ * @param[in] gamma Gamma values tensor info. 1 dimension with size equal to the feature maps [FM]. If not provided, default value for gamma is 1. Data types supported: Same as @p input
+ * @param[in] epsilon Small value to avoid division with zero. Default value is 0.001f.
+ * @param[in] conv_info Contains padding and stride information described in @ref PadStrideInfo.
+ * @param[in] num_groups Number of groups when performing a grouped convolution. num_groups != 1 is only supported for NCHW data layout
+ * @param[in] fast_math Enable fast math computation. In case this flag were set, the function could dispatch the fastest implementation
+ * available which may introduce a drop of accuracy as well. Default is false
+ * @param[in] fused_act Activation layer information in case of a fused activation.
+ *
+ */
+ void configure(TensorType *input,
+ TensorType *weights,
+ TensorType *bias,
+ TensorType *output,
+ const TensorType *mean,
+ const TensorType *var,
+ const TensorType *beta,
+ const TensorType *gamma,
+ float epsilon, const PadStrideInfo &conv_info, unsigned int num_groups, bool fast_math, ActivationLayerInfo const &fused_act)
+ {
+ // We don't run any validate, as we assume that the layers have been already validated
+ const bool has_bias = (bias != nullptr);
+ const TensorType *bias_to_use;
+
+ // We check if the layer has a bias. If yes, use it in-place. If not, we need to create one
+ // as batch normalization might end up with a bias != 0
+ if(has_bias)
+ {
+ _fused_batch_norm_layer.configure(weights, mean, var, nullptr, nullptr, bias, beta, gamma, epsilon);
+ bias_to_use = bias;
+ }
+ else
+ {
+ _fused_batch_norm_layer.configure(weights, mean, var, nullptr, &_fused_bias, nullptr, beta, gamma, epsilon);
+ bias_to_use = &_fused_bias;
+ }
+
+ _conv_layer.configure(input, weights, bias_to_use, output, conv_info, WeightsInfo(), Size2D(1U, 1U), fused_act, fast_math, num_groups);
+
+ if(!has_bias)
+ {
+ _fused_bias.allocator()->allocate();
+ }
+ }
+
+ // Inherited methods overridden:
+ void run()
+ {
+ prepare();
+ _conv_layer.run();
+ }
+
+ void prepare()
+ {
+ if(!_is_prepared)
+ {
+ _fused_batch_norm_layer.run();
+ _is_prepared = true;
+ }
+ }
+
+private:
+ typename FusedLayerTypes::ConvolutionLayer _conv_layer;
+ typename FusedLayerTypes::FuseBatchNormalization _fused_batch_norm_layer;
+ TensorConcreteType _fused_bias;
+ bool _is_prepared;
+};
+} // namespace backends
+} // namespace graph
+} // namespace arm_compute
+
+#endif /* __ARM_COMPUTE_GRAPH_BACKENDS_FUSED_CONVOLUTION_BATCH_NORMAZLIZATION_FUNCTION_H__ */
diff --git a/arm_compute/graph/mutators/NodeFusionMutator.h b/arm_compute/graph/mutators/NodeFusionMutator.h
index 8f16c65..b9ca464 100644
--- a/arm_compute/graph/mutators/NodeFusionMutator.h
+++ b/arm_compute/graph/mutators/NodeFusionMutator.h
@@ -1,5 +1,5 @@
/*
- * Copyright (c) 2018 ARM Limited.
+ * Copyright (c) 2018-2019 ARM Limited.
*
* SPDX-License-Identifier: MIT
*
@@ -24,21 +24,13 @@
#ifndef __ARM_COMPUTE_GRAPH_NODE_FUSION_MUTATOR_H__
#define __ARM_COMPUTE_GRAPH_NODE_FUSION_MUTATOR_H__
+#include "arm_compute/graph/Graph.h"
#include "arm_compute/graph/IGraphMutator.h"
namespace arm_compute
{
namespace graph
{
-namespace detail
-{
-/** Fused batch normalization with activation
- *
- * @param[in] g Graph to perform operation fusion on
- */
-void fuse_batch_norm_with_activation(Graph &g);
-} // namespace detail
-
/** Mutation pass to fuss nodes */
class NodeFusionMutator final : public IGraphMutator
{
diff --git a/arm_compute/graph/nodes/ActivationLayerNode.h b/arm_compute/graph/nodes/ActivationLayerNode.h
index 570351b..7231206 100644
--- a/arm_compute/graph/nodes/ActivationLayerNode.h
+++ b/arm_compute/graph/nodes/ActivationLayerNode.h
@@ -1,5 +1,5 @@
/*
- * Copyright (c) 2018 ARM Limited.
+ * Copyright (c) 2018-2019 ARM Limited.
*
* SPDX-License-Identifier: MIT
*
@@ -51,6 +51,9 @@
TensorDescriptor configure_output(size_t idx) const override;
void accept(INodeVisitor &v) override;
+public:
+ static constexpr NodeType node_type = NodeType::ActivationLayer;
+
private:
ActivationLayerInfo _info;
};
diff --git a/arm_compute/graph/nodes/FusedConvolutionBatchNormalizationNode.h b/arm_compute/graph/nodes/FusedConvolutionBatchNormalizationNode.h
new file mode 100644
index 0000000..9b0f5b7
--- /dev/null
+++ b/arm_compute/graph/nodes/FusedConvolutionBatchNormalizationNode.h
@@ -0,0 +1,144 @@
+/*
+ * Copyright (c) 2019 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#ifndef __ARM_COMPUTE_GRAPH_FUSED_CONVOLUTION_BATCH_NORMALIZATION_NODE_H__
+#define __ARM_COMPUTE_GRAPH_FUSED_CONVOLUTION_BATCH_NORMALIZATION_NODE_H__
+
+#include "arm_compute/graph/INode.h"
+
+namespace arm_compute
+{
+namespace graph
+{
+/** Batch Normalization node */
+class FusedConvolutionBatchNormalizationNode final : public INode
+{
+public:
+ /** Constructor
+ *
+ * @param[in] epsilon Epsilon parameter.
+ * @param[in] info Convolution layer attributes.
+ * @param[in] num_groups (Optional) Number of groups (Defaults to 1)
+ * @param[in] method (Optional) Convolution method to use
+ * @param[in] fast_math_hint (Optional) Fast math hint
+ * @param[in] out_quant_info (Optional) Output quantization info
+ * @param[in] fused_activation (Optional) Fused activation layer. Disabled if not specified
+ */
+ FusedConvolutionBatchNormalizationNode(float epsilon, PadStrideInfo info,
+ unsigned int num_groups = 1,
+ ConvolutionMethod method = ConvolutionMethod::Default,
+ FastMathHint fast_math_hint = FastMathHint::Disabled,
+ QuantizationInfo out_quant_info = QuantizationInfo(), ActivationLayerInfo fused_activation = ActivationLayerInfo());
+
+ /** Epsilon parameter accessor
+ *
+ * @return Epsilon parameter
+ */
+ float epsilon() 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);
+
+ /** Computes convolution output descriptor
+ *
+ * @param[in] input_descriptor Input descriptor
+ * @param[in] weights_descriptor Weights descriptor
+ * @param[in] info Convolution operation attributes
+ *
+ * @return Output descriptor
+ */
+ static TensorDescriptor compute_output_descriptor(const TensorDescriptor &input_descriptor,
+ const TensorDescriptor &weights_descriptor,
+ const PadStrideInfo &info);
+
+ /** Sets the convolution layer method to use
+ *
+ * @param[in] method Method to use for convolution
+ */
+ void set_convolution_method(ConvolutionMethod method);
+
+ /** Number of groups in convolution accessor
+ *
+ * @return Number of groups in convolution
+ */
+ unsigned int num_groups() const;
+
+ /** Convolution layer method accessor
+ *
+ * @note This is an indication on which convolution layer implementation to use,
+ * if it fails to be created the library's heuristic approach will be used
+ *
+ * @return Convolution layer method to be used by the node
+ */
+ ConvolutionMethod convolution_method() const;
+
+ /** Sets the fast math fast hint
+ *
+ * @param[in] hint Hint to use for convolution
+ */
+ void set_fast_math_hint(FastMathHint hint);
+
+ /** Fast math hint accessor
+ *
+ * @return Fast math hint to be used by the node
+ */
+ FastMathHint fast_math_hint() const;
+
+ /** Convolution metadata accessor
+ *
+ * @return Convolution information
+ */
+ PadStrideInfo convolution_info() const;
+
+ // Inherited overridden methods:
+ NodeType type() const override;
+ bool forward_descriptors() override;
+ TensorDescriptor configure_output(size_t idx) const override;
+ void accept(INodeVisitor &v) override;
+
+public:
+ static constexpr NodeType node_type = NodeType::FusedConvolutionBatchNormalizationLayer;
+
+private:
+ float _epsilon;
+
+ PadStrideInfo _info;
+ unsigned int _num_groups;
+ ConvolutionMethod _method;
+ FastMathHint _fast_math_hint;
+ QuantizationInfo _out_quant_info;
+ ActivationLayerInfo _fused_activation;
+};
+
+} // namespace graph
+} // namespace arm_compute
+#endif /* __ARM_COMPUTE_GRAPH_BATCH_NORMALIZATION_LAYER_NODE_H__ */
diff --git a/arm_compute/graph/nodes/Nodes.h b/arm_compute/graph/nodes/Nodes.h
index 2406485..e23b2b9 100644
--- a/arm_compute/graph/nodes/Nodes.h
+++ b/arm_compute/graph/nodes/Nodes.h
@@ -38,6 +38,7 @@
#include "arm_compute/graph/nodes/EltwiseLayerNode.h"
#include "arm_compute/graph/nodes/FlattenLayerNode.h"
#include "arm_compute/graph/nodes/FullyConnectedLayerNode.h"
+#include "arm_compute/graph/nodes/FusedConvolutionBatchNormalizationNode.h"
#include "arm_compute/graph/nodes/GenerateProposalsLayerNode.h"
#include "arm_compute/graph/nodes/InputNode.h"
#include "arm_compute/graph/nodes/NormalizationLayerNode.h"
diff --git a/arm_compute/graph/nodes/NodesFwd.h b/arm_compute/graph/nodes/NodesFwd.h
index cbda309..80576d4 100644
--- a/arm_compute/graph/nodes/NodesFwd.h
+++ b/arm_compute/graph/nodes/NodesFwd.h
@@ -44,6 +44,7 @@
class EltwiseLayerNode;
class FlattenLayerNode;
class FullyConnectedLayerNode;
+class FusedConvolutionBatchNormalizationNode;
class GenerateProposalsLayerNode;
class InputNode;
class NormalizationLayerNode;
diff --git a/arm_compute/graph/printers/DotGraphPrinter.h b/arm_compute/graph/printers/DotGraphPrinter.h
index d4cf692..9d2ea46 100644
--- a/arm_compute/graph/printers/DotGraphPrinter.h
+++ b/arm_compute/graph/printers/DotGraphPrinter.h
@@ -1,5 +1,5 @@
/*
- * Copyright (c) 2018 ARM Limited.
+ * Copyright (c) 2018-2019 ARM Limited.
*
* SPDX-License-Identifier: MIT
*
@@ -56,6 +56,7 @@
void visit(ConvolutionLayerNode &n) override;
void visit(DepthwiseConvolutionLayerNode &n) override;
void visit(EltwiseLayerNode &n) override;
+ void visit(FusedConvolutionBatchNormalizationNode &n) override;
void visit(NormalizationLayerNode &n) override;
void visit(PoolingLayerNode &n) override;
void default_visit() override;
diff --git a/src/core/CL/cl_kernels/batchnormalization_layer.cl b/src/core/CL/cl_kernels/batchnormalization_layer.cl
index dfd16e0..60307bc 100644
--- a/src/core/CL/cl_kernels/batchnormalization_layer.cl
+++ b/src/core/CL/cl_kernels/batchnormalization_layer.cl
@@ -1,5 +1,5 @@
/*
- * Copyright (c) 2017-2018 ARM Limited.
+ * Copyright (c) 2017-2019 ARM Limited.
*
* SPDX-License-Identifier: MIT
*
@@ -341,22 +341,10 @@
Vector bn_mean = CONVERT_TO_VECTOR_STRUCT_NO_STEP(bn_mean);
Vector bn_var = CONVERT_TO_VECTOR_STRUCT_NO_STEP(bn_var);
- // In-place ops
-#ifdef IN_PLACE_W
- Tensor4D fused_w = conv_w;
-#else /* IN_PLACE_W */
- Tensor4D fused_w = CONVERT_TO_TENSOR4D_STRUCT(fused_w, NUM_CHANNELS);
-#endif /* IN_PLACE */
-#ifdef IN_PLACE_B
- Vector fused_b = conv_b;
-#else /* IN_PLACE_W */
- Vector fused_b = CONVERT_TO_VECTOR_STRUCT_NO_STEP(fused_b);
-#endif /* IN_PLACE */
-
// Conditional ops
#ifdef HAS_BIAS
Vector conv_b = CONVERT_TO_VECTOR_STRUCT_NO_STEP(conv_b);
-#endif /* USE_DEFAULT_BETA */
+#endif /* HAS_BIAS */
#ifndef USE_DEFAULT_BETA
Vector bn_beta = CONVERT_TO_VECTOR_STRUCT_NO_STEP(bn_beta);
#endif /* USE_DEFAULT_BETA */
@@ -364,6 +352,19 @@
Vector bn_gamma = CONVERT_TO_VECTOR_STRUCT_NO_STEP(bn_gamma);
#endif /* USE_DEFAULT_GAMMA */
+ // In-place ops
+#ifdef IN_PLACE_W
+ Tensor4D fused_w = conv_w;
+ uint fused_w_stride_x = conv_w_stride_x;
+#else /* IN_PLACE_W */
+ Tensor4D fused_w = CONVERT_TO_TENSOR4D_STRUCT(fused_w, NUM_CHANNELS);
+#endif /* IN_PLACE_W */
+#ifdef IN_PLACE_B
+ Vector fused_b = conv_b;
+#else /* IN_PLACE_B */
+ Vector fused_b = CONVERT_TO_VECTOR_STRUCT_NO_STEP(fused_b);
+#endif /* IN_PLACE_B */
+
const int current_slice = get_global_id(2) / NUM_CHANNELS;
#if defined(VEC_SIZE) && defined(LAST_ACCESSED_X)
diff --git a/src/graph/backends/CL/CLFunctionsFactory.cpp b/src/graph/backends/CL/CLFunctionsFactory.cpp
index b9e3ddc..7473ff4 100644
--- a/src/graph/backends/CL/CLFunctionsFactory.cpp
+++ b/src/graph/backends/CL/CLFunctionsFactory.cpp
@@ -40,7 +40,8 @@
/** Target specific information structure used to pass information to the layer templates */
struct CLTargetInfo
{
- using TensorType = arm_compute::ICLTensor;
+ using TensorType = arm_compute::ICLTensor;
+ using TensorConcreteType = CLTensor;
static Target TargetType;
};
@@ -69,6 +70,14 @@
using Subtraction = CLArithmeticSubtraction;
using Multiplication = CLPixelWiseMultiplication;
};
+
+/** Function and tensor types to be used inside a CL fused convolution/batch normalization layer */
+struct CLFusedLayerTypes
+{
+ using ConvolutionLayer = CLConvolutionLayer;
+ using FuseBatchNormalization = CLFuseBatchNormalization;
+};
+
// TODO (isagot01): Remove once we support heterogeneous scheduling at function level
/** Wrapper for the CPP Function in the OpenCL backend **/
class CPPWrapperFunction : public IFunction
@@ -192,6 +201,8 @@
return detail::create_flatten_layer<CLFlattenLayer, CLTargetInfo>(*polymorphic_downcast<FlattenLayerNode *>(node));
case NodeType::FullyConnectedLayer:
return detail::create_fully_connected_layer<CLFullyConnectedLayer, CLTargetInfo>(*polymorphic_downcast<FullyConnectedLayerNode *>(node), ctx);
+ case NodeType::FusedConvolutionBatchNormalizationLayer:
+ return detail::create_fused_convolution_batch_normalization_layer<CLFusedLayerTypes, CLTargetInfo>(*polymorphic_downcast<FusedConvolutionBatchNormalizationNode *>(node));
case NodeType::GenerateProposalsLayer:
return detail::create_generate_proposals_layer<CLGenerateProposalsLayer, CLTargetInfo>(*polymorphic_downcast<GenerateProposalsLayerNode *>(node), ctx);
case NodeType::NormalizationLayer:
diff --git a/src/graph/backends/NEON/NEFunctionFactory.cpp b/src/graph/backends/NEON/NEFunctionFactory.cpp
index dc987dd..f23845c 100644
--- a/src/graph/backends/NEON/NEFunctionFactory.cpp
+++ b/src/graph/backends/NEON/NEFunctionFactory.cpp
@@ -1,5 +1,5 @@
/*
- * Copyright (c) 2018 ARM Limited.
+ * Copyright (c) 2018-2019 ARM Limited.
*
* SPDX-License-Identifier: MIT
*
@@ -46,7 +46,8 @@
/** Target specific information structure used to pass information to the layer templates */
struct NETargetInfo
{
- using TensorType = arm_compute::ITensor;
+ using TensorType = arm_compute::ITensor;
+ using TensorConcreteType = arm_compute::Tensor;
static Target TargetType;
};
@@ -76,6 +77,13 @@
using Multiplication = NEPixelWiseMultiplication;
};
+/** Function and tensor types to be used inside a NEON fused convolution/batch normalization layer */
+struct NEFusedLayerTypes
+{
+ using ConvolutionLayer = NEConvolutionLayer;
+ using FuseBatchNormalization = NEFuseBatchNormalization;
+};
+
namespace detail
{
// Specialized functions
@@ -210,6 +218,8 @@
return detail::create_flatten_layer<NEFlattenLayer, NETargetInfo>(*polymorphic_downcast<FlattenLayerNode *>(node));
case NodeType::FullyConnectedLayer:
return detail::create_fully_connected_layer<NEFullyConnectedLayer, NETargetInfo>(*polymorphic_downcast<FullyConnectedLayerNode *>(node), ctx);
+ case NodeType::FusedConvolutionBatchNormalizationLayer:
+ return detail::create_fused_convolution_batch_normalization_layer<NEFusedLayerTypes, NETargetInfo>(*polymorphic_downcast<FusedConvolutionBatchNormalizationNode *>(node));
case NodeType::NormalizationLayer:
return detail::create_normalization_layer<NENormalizationLayer, NETargetInfo>(*polymorphic_downcast<NormalizationLayerNode *>(node), ctx);
case NodeType::PermuteLayer:
diff --git a/src/graph/mutators/NodeFusionMutator.cpp b/src/graph/mutators/NodeFusionMutator.cpp
index 9dc02d1..445748c 100644
--- a/src/graph/mutators/NodeFusionMutator.cpp
+++ b/src/graph/mutators/NodeFusionMutator.cpp
@@ -1,5 +1,5 @@
/*
- * Copyright (c) 2018 ARM Limited.
+ * Copyright (c) 2018-2019 ARM Limited.
*
* SPDX-License-Identifier: MIT
*
@@ -23,9 +23,11 @@
*/
#include "arm_compute/graph/mutators/NodeFusionMutator.h"
-#include "arm_compute/graph/Graph.h"
+#include "arm_compute/graph/GraphBuilder.h"
#include "arm_compute/graph/Logger.h"
#include "arm_compute/graph/Utils.h"
+#include "arm_compute/graph/backends/BackendRegistry.h"
+#include "arm_compute/graph/nodes/FusedConvolutionBatchNormalizationNode.h"
#include "arm_compute/graph/nodes/Nodes.h"
#include "arm_compute/core/utils/misc/Cast.h"
@@ -38,69 +40,156 @@
{
namespace detail
{
+void fuse_convolution_with_batch_normalization(Graph &g, const Edge *output_edge)
+{
+ ARM_COMPUTE_ERROR_ON(output_edge == nullptr);
+
+ auto *conv_node = arm_compute::utils::cast::polymorphic_downcast<ConvolutionLayerNode *>(output_edge->producer());
+ auto *bn_node = arm_compute::utils::cast::polymorphic_downcast<BatchNormalizationLayerNode *>(output_edge->consumer());
+
+ // Not fusing if number of groups is greater than 1
+ if(conv_node->num_groups() > 1)
+ {
+ return;
+ }
+
+ ARM_COMPUTE_LOG_GRAPH_VERBOSE("Fusing convolution node with ID : " << output_edge->producer_id()
+ << " with BatchNormalization Layer node with ID : " << output_edge->consumer_id() << std::endl);
+
+ // Prevent fusion if fused node has an output accessor
+ if(conv_node->output(0)->accessor() == nullptr)
+ {
+ const Target assigned_target = conv_node->assigned_target();
+
+ // Extract conv inputs
+ const auto conv_input_id = conv_node->input_edge(0)->producer_id();
+ const auto conv_weights_id = conv_node->input_edge(1)->producer_id();
+ const auto out_quant_info = conv_node->output(0)->desc().quant_info;
+ const auto conv_info = conv_node->convolution_info();
+ const auto conv_method = conv_node->convolution_method();
+ const auto num_groups = conv_node->num_groups();
+ const auto act_info = bn_node->fused_activation();
+ FastMathHint fast_math_hint = conv_node->fast_math_hint();
+
+ // Extract bn inputs
+ const auto bn_mean_id = bn_node->input_edge(1)->producer_id();
+ const auto bn_var_id = bn_node->input_edge(2)->producer_id();
+ const auto bn_beta_id = bn_node->input_edge(3)->producer_id();
+ const auto bn_gamma_id = bn_node->input_edge(4)->producer_id();
+ const auto epsilon = bn_node->epsilon();
+
+ // Create the fused node
+ const NodeID fused_id = g.add_node<FusedConvolutionBatchNormalizationNode>(epsilon, conv_info, num_groups, conv_method, fast_math_hint, out_quant_info, act_info);
+
+ if(conv_node->input_edge(2) != nullptr)
+ {
+ auto conv_bias_id = conv_node->input_edge(2)->producer_id();
+ g.add_connection(conv_bias_id, 0, fused_id, 2);
+ }
+
+ // Add connections from the conv/batch_norm inputs to the fused node
+ g.add_connection(conv_input_id, 0, fused_id, 0);
+ g.add_connection(conv_weights_id, 0, fused_id, 1);
+ g.add_connection(bn_mean_id, 0, fused_id, 3);
+ g.add_connection(bn_var_id, 0, fused_id, 4);
+ g.add_connection(bn_beta_id, 0, fused_id, 5);
+ g.add_connection(bn_gamma_id, 0, fused_id, 6);
+
+ auto fused_node = g.node(fused_id);
+ std::vector<NodeIdxPair> bn_driving_nodes = get_driving_nodes(*bn_node);
+
+ // Extract batch normalization node accessor if any
+ auto bn_node_accessor = bn_node->output(0)->extract_accessor();
+ auto bn_node_name = bn_node->name();
+
+ // Remove batch normalization node
+ g.remove_node(bn_node->id());
+
+ // Get driving nodes of batch normalization node
+ for(auto &driving_node : bn_driving_nodes)
+ {
+ g.add_connection(fused_id, 0, driving_node.node_id, driving_node.index);
+ configure_tensor(fused_node->output(0));
+ }
+ // Update fused node outputs
+ fused_node->output(0)->set_accessor(std::move(bn_node_accessor));
+ fused_node->set_assigned_target(assigned_target);
+ fused_node->set_common_node_parameters(NodeParams{ conv_node->name() + "+" + bn_node_name, assigned_target });
+
+ // Remove convolution node
+ g.remove_node(conv_node->id());
+ }
+ else
+ {
+ ARM_COMPUTE_LOG_GRAPH_VERBOSE("Prevented fusion of convolution with batch normalization due to the presence of an output accessor\n");
+ }
+}
+
template <typename N>
-void fuse_node_with_activation(Graph &g,
- const std::set<Activation> &supported_fused_activations,
- std::function<bool(INode &)> const &prec)
+void fuse_node_with_activation(Graph &g, const Edge *output_edge, const std::set<Activation> &supported_fused_activations)
+{
+ ARM_COMPUTE_ERROR_ON(output_edge == nullptr);
+
+ auto *n_node = arm_compute::utils::cast::polymorphic_downcast<N *>(output_edge->producer());
+ auto *act_node = arm_compute::utils::cast::polymorphic_downcast<ActivationLayerNode *>(output_edge->consumer());
+
+ ARM_COMPUTE_ERROR_ON(act_node->output(0) == nullptr || n_node->output(0) == nullptr);
+
+ // Check if activation is supported for fusion
+ if(supported_fused_activations.count(act_node->activation_info().activation()) == 0)
+ {
+ return;
+ }
+
+ ARM_COMPUTE_LOG_GRAPH_VERBOSE("Fusing node with ID : " << output_edge->producer_id()
+ << " with Activation Layer node with ID : " << output_edge->consumer_id() << std::endl);
+
+ // Prevent fusion if fused node has an output accessor
+ if(n_node->output(0)->accessor() == nullptr)
+ {
+ // Get driving nodes of activation node
+ std::vector<NodeIdxPair> act_driving_nodes = get_driving_nodes(*act_node);
+
+ // Set activation info to fused node
+ n_node->set_fused_activation(act_node->activation_info());
+
+ // Extract activation node accessor if any
+ auto act_node_accessor = act_node->output(0)->extract_accessor();
+
+ // Remove activation node
+ g.remove_node(act_node->id());
+
+ // Update fused node outputs
+ for(auto &driving_node : act_driving_nodes)
+ {
+ g.add_connection(n_node->id(), 0, driving_node.node_id, driving_node.index);
+ }
+
+ // Update accessor to fused node
+ n_node->output(0)->set_accessor(std::move(act_node_accessor));
+ }
+ else
+ {
+ ARM_COMPUTE_LOG_GRAPH_VERBOSE("Prevented fusion of node with activation due to the presence of an output accessor\n");
+ }
+}
+
+template <typename N1, typename N2, typename F, typename... Args>
+void fuse_layer(Graph &g, std::function<bool(INode &)> const &prec, const F fuse_fcn, Args &&... optional_arguments)
{
// Not interested in the order of nodes
for(auto &node : g.nodes())
{
// Check if the node is of type N and not a branching node
- if(node && node->type() == N::node_type && node->output_edges().size() == 1)
+ if(node && node->type() == N1::node_type && node->output_edges().size() == 1)
{
- auto output_edge_id = *node->output_edges().begin();
- auto output_edge = g.edge(output_edge_id);
+ const auto output_edge_id = *node->output_edges().begin();
+ const auto output_edge = g.edge(output_edge_id);
+
// Check if following node is an activation layer node
- if((output_edge != nullptr) && (output_edge->consumer() != nullptr) && (output_edge->consumer()->type() == NodeType::ActivationLayer))
+ if((output_edge != nullptr) && (output_edge->consumer() != nullptr) && (output_edge->consumer()->type() == N2::node_type) && prec(*output_edge->producer()))
{
- auto *n_node = arm_compute::utils::cast::polymorphic_downcast<N *>(output_edge->producer());
- auto *act_node = arm_compute::utils::cast::polymorphic_downcast<ActivationLayerNode *>(output_edge->consumer());
-
- ARM_COMPUTE_ERROR_ON(act_node->output(0) == nullptr || n_node->output(0) == nullptr);
-
- // Check given precondition
- if(!prec(*n_node))
- {
- continue;
- }
- // Check if activation is supported for fusion
- if(supported_fused_activations.count(act_node->activation_info().activation()) == 0)
- {
- continue;
- }
-
- ARM_COMPUTE_LOG_GRAPH_VERBOSE("Fusing node with ID : " << output_edge->producer_id()
- << " with Activation Layer node with ID : " << output_edge->consumer_id() << std::endl);
-
- // Prevent fusion if fused node has an output accessor
- if(n_node->output(0)->accessor() == nullptr)
- {
- // Get driving nodes of activation node
- std::vector<NodeIdxPair> act_driving_nodes = get_driving_nodes(*act_node);
-
- // Set activation info to fused node
- n_node->set_fused_activation(act_node->activation_info());
-
- // Extract activation node accessor if any
- auto act_node_accessor = act_node->output(0)->extract_accessor();
-
- // Remove activation node
- g.remove_node(act_node->id());
-
- // Update fused node outputs
- for(auto &driving_node : act_driving_nodes)
- {
- g.add_connection(n_node->id(), 0, driving_node.node_id, driving_node.index);
- }
-
- // Update accessor to fused node
- n_node->output(0)->set_accessor(std::move(act_node_accessor));
- }
- else
- {
- ARM_COMPUTE_LOG_GRAPH_VERBOSE("Prevented fusion of node with activation due to the presence of an output accessor\n");
- }
+ fuse_fcn(g, output_edge, optional_arguments...);
}
}
}
@@ -129,9 +218,10 @@
};
// Fusion mutations
- detail::fuse_node_with_activation<BatchNormalizationLayerNode>(g, supported_fused_activations, empty_prec);
- detail::fuse_node_with_activation<ConvolutionLayerNode>(g, supported_fused_activations, empty_prec);
- detail::fuse_node_with_activation<DepthwiseConvolutionLayerNode>(g, supported_fused_activations, qs8_prec);
+ detail::fuse_layer<BatchNormalizationLayerNode, ActivationLayerNode>(g, empty_prec, detail::fuse_node_with_activation<BatchNormalizationLayerNode>, supported_fused_activations);
+ detail::fuse_layer<ConvolutionLayerNode, ActivationLayerNode>(g, empty_prec, detail::fuse_node_with_activation<ConvolutionLayerNode>, supported_fused_activations);
+ detail::fuse_layer<DepthwiseConvolutionLayerNode, ActivationLayerNode>(g, qs8_prec, detail::fuse_node_with_activation<DepthwiseConvolutionLayerNode>, supported_fused_activations);
+ detail::fuse_layer<ConvolutionLayerNode, BatchNormalizationLayerNode>(g, empty_prec, detail::fuse_convolution_with_batch_normalization);
}
} // namespace graph
} // namespace arm_compute
diff --git a/src/graph/nodes/ActivationLayerNode.cpp b/src/graph/nodes/ActivationLayerNode.cpp
index 414684c..85cb10b 100644
--- a/src/graph/nodes/ActivationLayerNode.cpp
+++ b/src/graph/nodes/ActivationLayerNode.cpp
@@ -1,5 +1,5 @@
/*
- * Copyright (c) 2018 ARM Limited.
+ * Copyright (c) 2018-2019 ARM Limited.
*
* SPDX-License-Identifier: MIT
*
@@ -67,7 +67,7 @@
NodeType ActivationLayerNode::type() const
{
- return NodeType::ActivationLayer;
+ return ActivationLayerNode::node_type;
}
void ActivationLayerNode::accept(INodeVisitor &v)
@@ -75,4 +75,4 @@
v.visit(*this);
}
} // namespace graph
-} // namespace arm_compute
\ No newline at end of file
+} // namespace arm_compute
diff --git a/src/graph/nodes/FusedConvolutionBatchNormalizationNode.cpp b/src/graph/nodes/FusedConvolutionBatchNormalizationNode.cpp
new file mode 100644
index 0000000..27a348f
--- /dev/null
+++ b/src/graph/nodes/FusedConvolutionBatchNormalizationNode.cpp
@@ -0,0 +1,152 @@
+/*
+ * Copyright (c) 2019 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#include "arm_compute/graph/nodes/FusedConvolutionBatchNormalizationNode.h"
+
+#include "arm_compute/core/Utils.h"
+#include "arm_compute/graph/Graph.h"
+#include "arm_compute/graph/INodeVisitor.h"
+#include "arm_compute/graph/Utils.h"
+
+namespace arm_compute
+{
+namespace graph
+{
+FusedConvolutionBatchNormalizationNode::FusedConvolutionBatchNormalizationNode(float epsilon, PadStrideInfo info,
+ unsigned int num_groups,
+ ConvolutionMethod method,
+ FastMathHint fast_math_hint,
+ QuantizationInfo out_quant_info, ActivationLayerInfo fused_activation)
+ : _epsilon(epsilon), _info(std::move(info)), _num_groups(num_groups), _method(method), _fast_math_hint(fast_math_hint), _out_quant_info(out_quant_info), _fused_activation(fused_activation)
+{
+ _input_edges.resize(7, EmptyEdgeID);
+ _outputs.resize(1, NullTensorID);
+}
+
+void FusedConvolutionBatchNormalizationNode::set_convolution_method(ConvolutionMethod method)
+{
+ _method = method;
+}
+
+float FusedConvolutionBatchNormalizationNode::epsilon() const
+{
+ return _epsilon;
+}
+
+ConvolutionMethod FusedConvolutionBatchNormalizationNode::convolution_method() const
+{
+ return _method;
+}
+
+void FusedConvolutionBatchNormalizationNode::set_fast_math_hint(FastMathHint hint)
+{
+ _fast_math_hint = hint;
+}
+
+FastMathHint FusedConvolutionBatchNormalizationNode::fast_math_hint() const
+{
+ return _fast_math_hint;
+}
+
+PadStrideInfo FusedConvolutionBatchNormalizationNode::convolution_info() const
+{
+ return _info;
+}
+
+unsigned int FusedConvolutionBatchNormalizationNode::num_groups() const
+{
+ return _num_groups;
+}
+
+ActivationLayerInfo FusedConvolutionBatchNormalizationNode::fused_activation() const
+{
+ return _fused_activation;
+}
+
+void FusedConvolutionBatchNormalizationNode::set_fused_activation(ActivationLayerInfo fused_activation)
+{
+ _fused_activation = fused_activation;
+}
+
+TensorDescriptor FusedConvolutionBatchNormalizationNode::compute_output_descriptor(const TensorDescriptor &input_descriptor,
+ const TensorDescriptor &weights_descriptor,
+ const PadStrideInfo &info)
+{
+ unsigned int output_width = 0;
+ unsigned int output_height = 0;
+
+ const unsigned int input_width = get_dimension_size(input_descriptor, DataLayoutDimension::WIDTH);
+ const unsigned int input_height = get_dimension_size(input_descriptor, DataLayoutDimension::HEIGHT);
+ const unsigned int kernel_width = get_dimension_size(weights_descriptor, DataLayoutDimension::WIDTH);
+ const unsigned int kernel_height = get_dimension_size(weights_descriptor, DataLayoutDimension::HEIGHT);
+
+ std::tie(output_width, output_height) = scaled_dimensions(input_width, input_height, kernel_width, kernel_height, info);
+
+ TensorDescriptor output_descriptor = input_descriptor;
+ output_descriptor.shape.set(get_dimension_idx(output_descriptor, DataLayoutDimension::WIDTH), output_width);
+ output_descriptor.shape.set(get_dimension_idx(output_descriptor, DataLayoutDimension::HEIGHT), output_height);
+ output_descriptor.shape.set(get_dimension_idx(output_descriptor, DataLayoutDimension::CHANNEL), weights_descriptor.shape[3]);
+
+ return output_descriptor;
+}
+
+bool FusedConvolutionBatchNormalizationNode::forward_descriptors()
+{
+ if((input_id(0) != NullTensorID) && (input_id(1) != NullTensorID) && (output_id(0) != NullTensorID))
+ {
+ Tensor *dst = output(0);
+ ARM_COMPUTE_ERROR_ON(dst == nullptr);
+ dst->desc() = configure_output(0);
+ return true;
+ }
+ return false;
+}
+
+TensorDescriptor FusedConvolutionBatchNormalizationNode::configure_output(size_t idx) const
+{
+ ARM_COMPUTE_UNUSED(idx);
+ const Tensor *src = input(0);
+ const Tensor *weights = input(1);
+
+ ARM_COMPUTE_ERROR_ON(src == nullptr || weights == nullptr);
+
+ TensorDescriptor output_info = compute_output_descriptor(src->desc(), weights->desc(), _info);
+ if(!_out_quant_info.empty())
+ {
+ output_info.quant_info = _out_quant_info;
+ }
+
+ return output_info;
+}
+
+NodeType FusedConvolutionBatchNormalizationNode::type() const
+{
+ return FusedConvolutionBatchNormalizationNode::node_type;
+}
+
+void FusedConvolutionBatchNormalizationNode::accept(INodeVisitor &v)
+{
+ v.visit(*this);
+}
+} // namespace graph
+} // namespace arm_compute
diff --git a/src/graph/printers/DotGraphPrinter.cpp b/src/graph/printers/DotGraphPrinter.cpp
index ef156ea..c939de1 100644
--- a/src/graph/printers/DotGraphPrinter.cpp
+++ b/src/graph/printers/DotGraphPrinter.cpp
@@ -1,5 +1,5 @@
/*
- * Copyright (c) 2018 ARM Limited.
+ * Copyright (c) 2018-2019 ARM Limited.
*
* SPDX-License-Identifier: MIT
*
@@ -77,6 +77,14 @@
_info = ss.str();
}
+void DotGraphVisitor::visit(FusedConvolutionBatchNormalizationNode &n)
+{
+ ARM_COMPUTE_UNUSED(n);
+ std::stringstream ss;
+ ss << "FusedConvolutionBatchNormalizationNode";
+ _info = ss.str();
+}
+
void DotGraphVisitor::visit(NormalizationLayerNode &n)
{
std::stringstream ss;