Changed location of extensibility guide (#10433)
* Changed location of extensibility guide * Removed hardware kernels legacy documentation * Changed all extension guild to new API * Removed Custom_Layers_Guide * Fixed build * Fixed some moments * Update docs/Extensibility_UG/Intro.md * Fixed build * Added more examples * Fixed typo * Fixed comments * Extend library topic * Fixed typo
This commit is contained in:
Ilya Churaev authored and 
GitHub committed
1 parent
a2f9963045
commit
1c18733ade
30 files changed
+268
-1632
No files matched your search
@@ -0,0 +1,114 @@
|
||||
# OpenVINO Extensibility Mechanism {#openvino_docs_Extensibility_UG_Intro}
|
||||
|
||||
@sphinxdirective
|
||||
|
||||
.. toctree::
|
||||
:maxdepth: 1
|
||||
:hidden:
|
||||
|
||||
openvino_docs_Extensibility_UG_add_openvino_ops
|
||||
|
||||
@endsphinxdirective
|
||||
|
||||
The Intel® Distribution of OpenVINO™ toolkit supports neural network models trained with multiple frameworks including
|
||||
TensorFlow, Caffe, MXNet, Kaldi, PaddlePaddle, and ONNX. The list of supported operations (layers) is different for
|
||||
each of the supported frameworks. To see the operations supported by your framework, refer to
|
||||
[Supported Framework Operations](../MO_DG/prepare_model/Supported_Frameworks_Layers.md).
|
||||
|
||||
Custom operations, that is those not included in the list, are not recognized by OpenVINO™ out-of-the-box. Therefore, creating Intermediate Representation (IR) for a model using them requires additional steps. This guide illustrates the workflow for running inference on topologies featuring custom operations, allowing you to plug in your own implementation for existing or completely new operations.
|
||||
|
||||
If your model contains operations not normally supported by OpenVINO™, the OpenVINO™ Extensibility API lets you add support for those custom operations and use one implementation for Model Optimizer and OpenVINO™ Runtime.
|
||||
|
||||
There are two steps to support inference of a model with custom operation(s):
|
||||
1. Add support for a [custom operation in the Model Optimizer](../MO_DG/prepare_model/customize_model_optimizer/Customize_Model_Optimizer.md) so
|
||||
the Model Optimizer can generate the IR with the operation.
|
||||
2. Create a custom operation in it as described in the [Custom Operation](add_openvino_ops.md).
|
||||
|
||||
## OpenVINO™ Extensions
|
||||
|
||||
An OpenVINO™ provides extensions for:
|
||||
|
||||
* [Custom OpenVINO™ Operation](add_openvino_ops.md):
|
||||
- Enables the creation of unsupported operations
|
||||
- Enables the use of `ov::Core::read_model` to read models with unsupported operations
|
||||
- Provides a shape inference mechanism for custom operations
|
||||
- Provides an evaluate method which allow to support the operation on CPU or perform constant folding
|
||||
|
||||
> **NOTE**: This documentation is written based on the [Template extension](https://github.com/openvinotoolkit/openvino/tree/master/docs/template_extension/new), which demonstrates extension development details. You can review the complete code, which is fully compilable and up-to-date, to see how it works.
|
||||
|
||||
## Load extensions to OpenVINO™ Runtime
|
||||
|
||||
To load the extensions to the `ov::Core` object, use the `ov::Core::add_extension` method, this method allows to load library with extensions or extensions from the code.
|
||||
|
||||
### Load extensions to core
|
||||
|
||||
Extensions can be loaded from code with `ov::Core::add_extension` method:
|
||||
|
||||
@sphinxdirective
|
||||
|
||||
.. tab:: C++
|
||||
|
||||
.. doxygensnippet:: docs/snippets/ov_extensions.cpp
|
||||
:language: cpp
|
||||
:fragment: add_extension
|
||||
|
||||
.. tab:: Python
|
||||
|
||||
.. doxygensnippet:: docs/snippets/ov_extensions.py
|
||||
:language: python
|
||||
:fragment: add_extension
|
||||
|
||||
@endsphinxdirective
|
||||
|
||||
### Create library with extensions
|
||||
|
||||
You need to create extension library in following cases:
|
||||
- Load extensions to Model Optimizer
|
||||
- Load extensions to Python application
|
||||
|
||||
If you want to create an extension library, for example in order to load these extensions to the Model Optimizer, you need to do next steps:
|
||||
Create an entry point for extension library. OpenVINO™ provides an `OPENVINO_CREATE_EXTENSIONS()` macro, which allows to define an entry point to a library with OpenVINO™ Extensions.
|
||||
This macro should have a vector of all OpenVINO™ Extensions as an argument.
|
||||
|
||||
Based on that, the declaration of an extension class can look as follows:
|
||||
|
||||
@snippet template_extension/new/ov_extension.cpp ov_extension:entry_point
|
||||
|
||||
To configure the build of your extension library, use the following CMake script:
|
||||
|
||||
@snippet template_extension/new/CMakeLists.txt cmake:extension
|
||||
|
||||
This CMake script finds the OpenVINO™ using the `find_package` CMake command.
|
||||
|
||||
To build the extension library, run the commands below:
|
||||
|
||||
```sh
|
||||
$ cd docs/template_extension/new
|
||||
$ mkdir build
|
||||
$ cd build
|
||||
$ cmake -DOpenVINO_DIR=<OpenVINO_DIR> ../
|
||||
$ cmake --build .
|
||||
```
|
||||
|
||||
After the build you can use path to your extension library to load your extensions to OpenVINO™ Runtime:
|
||||
|
||||
@sphinxdirective
|
||||
|
||||
.. tab:: C++
|
||||
|
||||
.. doxygensnippet:: docs/snippets/ov_extensions.cpp
|
||||
:language: cpp
|
||||
:fragment: add_extension_lib
|
||||
|
||||
.. tab:: Python
|
||||
|
||||
.. doxygensnippet:: docs/snippets/ov_extensions.py
|
||||
:language: python
|
||||
:fragment: add_extension_lib
|
||||
|
||||
@endsphinxdirective
|
||||
|
||||
## See Also
|
||||
|
||||
* [Using Inference Engine Samples](../OV_Runtime_UG/Samples_Overview.md)
|
||||
* [Hello Shape Infer SSD sample](../../samples/cpp/hello_reshape_ssd/README.md)
|
||||
@@ -0,0 +1,62 @@
|
||||
# Custom OpenVINO™ Operations {#openvino_docs_Extensibility_UG_add_openvino_ops}
|
||||
|
||||
OpenVINO™ Extension API allows you to register custom operations to support models with operations which OpenVINO™ does not support out-of-the-box.
|
||||
|
||||
## Operation Class
|
||||
|
||||
To add your custom operation, create a new class that extends `ov::Op`, which is in turn derived from `ov::Node`, the base class for all graph operations in OpenVINO™. To add `ov::Op` please include next file:
|
||||
|
||||
@snippet template_extension/new/identity.hpp op:common_include
|
||||
|
||||
Follow the steps below to add a custom operation:
|
||||
|
||||
1. Add the `OPENVINO_OP` macro which defines a `NodeTypeInfo` object that identifies the type of the operation to the graph users and helps with dynamic type resolution. The type info of an operation currently consists of a string operation identifier and a string for operation version.
|
||||
|
||||
2. Implement default constructor and constructors that optionally take the operation inputs and attributes as parameters.
|
||||
|
||||
3. Override the shape inference method `validate_and_infer_types`. This method is called multiple times during graph manipulations to determine the shapes and element types of the operations outputs. To access the input shapes and input element types, use the `get_input_partial_shape()` and `get_input_element_type()` methods of `ov::Node`. Set the inferred shape and element type of the output using `set_output_type`.
|
||||
|
||||
4. Override the `clone_with_new_inputs` method, which enables graph manipulation routines to create copies of this operation and connect it to different nodes during optimization.
|
||||
|
||||
5. Override the `visit_attributes` method, which enables serialization and deserialization of operation attributes. An `AttributeVisitor` is passed to the method, and the implementation is expected to walk over all the attributes in the op using the type-aware `on_attribute` helper. Helpers are already implemented for standard C++ types like `int64_t`, `float`, `bool`, `vector`, and for existing OpenVINO defined types.
|
||||
|
||||
6. Override `evaluate`, which is an optional method that enables fallback of some devices to this implementation and the application of constant folding if there is a custom operation on the constant branch. If your operation contains `evaluate` method you also need to override the `has_evaluate` method, this method allow to get information about availability of `evaluate` method for the operation.
|
||||
|
||||
7. Add the `OPENVINO_FRAMEWORK_MAP` macro if you want to map custom operation to framework operation with the same name. It is an optional macro which can be used for one to one mapping. In order to use this macro please include frontend specific headers:
|
||||
@snippet template_extension/new/identity.hpp op:frontend_include
|
||||
|
||||
Based on that, declaration of an operation class can look as follows:
|
||||
|
||||
@snippet template_extension/new/identity.hpp op:header
|
||||
|
||||
### Operation Constructors
|
||||
|
||||
OpenVINO™ operation contains two constructors:
|
||||
* Default constructor, which enables you to create an operation without attributes
|
||||
* Constructor that creates and validates an operation with specified inputs and attributes
|
||||
|
||||
@snippet template_extension/new/identity.cpp op:ctor
|
||||
|
||||
### `validate_and_infer_types()`
|
||||
|
||||
`ov::Node::validate_and_infer_types` method validates operation attributes and calculates output shapes using attributes of the operation.
|
||||
|
||||
@snippet template_extension/new/identity.cpp op:validate
|
||||
|
||||
### `clone_with_new_inputs()`
|
||||
|
||||
`ov::Node::clone_with_new_inputs` method creates a copy of the operation with new inputs.
|
||||
|
||||
@snippet template_extension/new/identity.cpp op:copy
|
||||
|
||||
### `visit_attributes()`
|
||||
|
||||
`ov::Node::visit_attributes` method enables you to visit all operation attributes.
|
||||
|
||||
@snippet template_extension/new/identity.cpp op:visit_attributes
|
||||
|
||||
### `evaluate()` and `has_evaluate()`
|
||||
|
||||
`ov::Node::evaluate` method enables you to apply constant folding to an operation.
|
||||
|
||||
@snippet template_extension/new/identity.cpp op:evaluate
|
||||
@@ -1,350 +0,0 @@
|
||||
# Custom Operations Guide {#openvino_docs_HOWTO_Custom_Layers_Guide}
|
||||
|
||||
The Intel® Distribution of OpenVINO™ toolkit supports neural network models trained with multiple frameworks, including
|
||||
TensorFlow, Caffe, MXNet, Kaldi, PaddlePaddle, and ONNX. The list of supported operations (layers) is different for
|
||||
each of the supported frameworks. To see the operations supported by your framework, refer to
|
||||
[Supported Framework Layers](../MO_DG/prepare_model/Supported_Frameworks_Layers.md).
|
||||
|
||||
Custom operations, that is those not included in the list, are not recognized by Model Optimizer out-of-the-box. Therefore, creating Intermediate Representation (IR) for a model using them requires additional steps. This guide illustrates the workflow for running inference on topologies featuring custom operations, allowing you to plug in your own implementation for existing or completely new operations.
|
||||
|
||||
> **NOTE**: *Layer* is a legacy term for *operation* which came from Caffe framework. Currently it is not used.
|
||||
> Refer to the [Deep Learning Network Intermediate Representation and Operation Sets in OpenVINO™](../MO_DG/IR_and_opsets.md)
|
||||
> for more information on the topic.
|
||||
|
||||
## Terms Used in This Guide
|
||||
|
||||
- *Intermediate Representation (IR)* — OpenVINO's Neural Network format used by Inference Engine. It abstracts different frameworks and describes model topology, operations parameters, and weights.
|
||||
|
||||
- *Operation* — an abstract concept of a math function selected for a specific purpose. Operations supported by
|
||||
OpenVINO™ are listed in the supported operation set provided in the [Available Operations Sets](../ops/opset.md).
|
||||
Examples of the operations are: [ReLU](../ops/activation/ReLU_1.md), [Convolution](../ops/convolution/Convolution_1.md),
|
||||
[Add](../ops/arithmetic/Add_1.md), etc.
|
||||
|
||||
- *Kernel* — The implementation of an operation function in the OpenVINO™ plugin, in this case, the math programmed (in
|
||||
C++ and OpenCL) to perform the operation for a target hardware (CPU or GPU).
|
||||
|
||||
- *Inference Engine Extension* — Device-specific module implementing custom operations (a set of kernels).
|
||||
|
||||
## Custom Operation Support Overview
|
||||
|
||||
There are three steps to support inference of a model with custom operation(s):
|
||||
1. Add support for a custom operation in [Model Optimizer](../MO_DG/Deep_Learning_Model_Optimizer_DevGuide.md) so
|
||||
that it can generate the IR with the operation.
|
||||
2. Create an operation set and implement a custom nGraph operation in it as described in the
|
||||
[Custom nGraph Operation](../OV_Runtime_UG/Extensibility_DG/AddingNGraphOps.md).
|
||||
3. Implement a customer operation in one of the [OpenVINO™ Runtime](../OV_Runtime_UG/openvino_intro.md)
|
||||
plugins to support inference of this operation using a particular target hardware (CPU, GPU or VPU).
|
||||
|
||||
To see the operations that are supported by each device plugin for the Inference Engine, refer to the
|
||||
[Supported Devices](../OV_Runtime_UG/supported_plugins/Supported_Devices.md).
|
||||
|
||||
> **NOTE**: If a device doesn't support a particular operation, an alternative to creating a new operation is to target
|
||||
> an additional device using the HETERO device. The [Heterogeneous execution](../OV_Runtime_UG/hetero_execution.md) may be
|
||||
> used to run an inference model on multiple devices allowing the unsupported operations on one device to "fallback" to
|
||||
> run on another device (e.g., CPU) that does support those operations.
|
||||
|
||||
### Custom Operation Support for the Model Optimizer
|
||||
|
||||
Model Optimizer model conversion pipeline is described in detail in "Model Conversion Pipeline" section of [Model Optimizer Extensibility](../MO_DG/prepare_model/customize_model_optimizer/Customize_Model_Optimizer.md). It is best to read that article first for a better understanding of the following material.
|
||||
|
||||
Model Optimizer provides an extensions mechanism to support new operations and implement custom model transformations to generate optimized IR. This mechanism is described in the "Model Optimizer Extensions" section of
|
||||
[Model Optimizer Extensibility](../MO_DG/prepare_model/customize_model_optimizer/Customize_Model_Optimizer.md).
|
||||
|
||||
Two types of Model Optimizer extensions should be implemented to support custom operations, at a minimum:
|
||||
1. Operation class for a new operation. This class stores information about the operation, its attributes, shape inference function, attributes to be saved to an IR and some others internally used attributes. Refer to the "Model Optimizer Operation" section of [Model Optimizer Extensibility](../MO_DG/prepare_model/customize_model_optimizer/Customize_Model_Optimizer.md) for detailed instructions on how to implement it.
|
||||
2. Operation attributes extractor. The extractor is responsible for parsing framework-specific representation of the
|
||||
operation and uses corresponding operation class to update graph node attributes with necessary attributes of the
|
||||
operation. Refer to the "Operation Extractor" section of
|
||||
[Model Optimizer Extensibility](../MO_DG/prepare_model/customize_model_optimizer/Customize_Model_Optimizer.md) for detailed instructions on how to implement it.
|
||||
|
||||
> **NOTE**: In some cases you may need to implement some transformation to support the operation. This topic is covered in the "Graph Transformation Extensions" section of [Model Optimizer Extensibility](../MO_DG/prepare_model/customize_model_optimizer/Customize_Model_Optimizer.md).
|
||||
|
||||
## Custom Operation Extensions for the Inference Engine
|
||||
|
||||
Inference Engine provides an extension mechanism to support new operations. This mechanism is described in [Inference Engine Extensibility Mechanism](../OV_Runtime_UG/Extensibility_DG/Intro.md).
|
||||
|
||||
Each device plugin includes a library of optimized implementations to execute known operations which must be extended to execute a custom operation. The custom operation extension is implemented according to the target device:
|
||||
|
||||
- Custom Operation CPU Extension
|
||||
- A compiled shared library (`.so` or `.dll`) needed by the CPU Plugin for executing the custom operation
|
||||
on a CPU. Refer to the [How to Implement Custom CPU Operations](../OV_Runtime_UG/Extensibility_DG/CPU_Kernel.md) for more
|
||||
details.
|
||||
- Custom Operation GPU Extension
|
||||
- OpenCL source code (.cl) for the custom operation kernel that will be compiled to execute on the GPU along with an operation description file (.xml) needed by the GPU Plugin for the custom operation kernel. Refer to the [How to Implement Custom GPU Operations](../OV_Runtime_UG/Extensibility_DG/GPU_Kernel.md) for more details.
|
||||
- Custom Operation VPU Extension
|
||||
- OpenCL source code (.cl) for the custom operation kernel that will be compiled to execute on the VPU along with an operation description file (.xml) needed by the VPU Plugin for the custom operation kernel. Refer to [How to Implement Custom Operations for VPU](../OV_Runtime_UG/Extensibility_DG/VPU_Kernel.md) for more details.
|
||||
|
||||
Also, it is necessary to implement nGraph custom operation according to [Custom nGraph Operation](../OV_Runtime_UG/Extensibility_DG/AddingNGraphOps.md) so the Inference Engine can read an IR with this
|
||||
operation and correctly infer output tensor shape and type.
|
||||
|
||||
## Enabling Magnetic Resonance Image Reconstruction Model
|
||||
This chapter provides step-by-step instructions on how to enable the magnetic resonance image reconstruction model implemented in the [repository](https://github.com/rmsouza01/Hybrid-CS-Model-MRI/) using a custom operation on CPU. The example is prepared for a model generated from the repository with hash `2ede2f96161ce70dcdc922371fe6b6b254aafcc8`.
|
||||
|
||||
### Download and Convert the Model to a Frozen TensorFlow Model Format
|
||||
The original pre-trained model is provided in the hdf5 format which is not supported by OpenVINO directly and needs to be converted to TensorFlow frozen model format first.
|
||||
|
||||
1. Download repository `https://github.com/rmsouza01/Hybrid-CS-Model-MRI`:<br>
|
||||
```bash
|
||||
git clone https://github.com/rmsouza01/Hybrid-CS-Model-MRI
|
||||
git checkout 2ede2f96161ce70dcdc922371fe6b6b254aafcc8
|
||||
```
|
||||
|
||||
2. Convert pre-trained `.hdf5` to a frozen `.pb` graph using the following script (tested with TensorFlow==1.15.0 and
|
||||
Keras==2.2.4) which should be executed from the root of the cloned repository:<br>
|
||||
```py
|
||||
import keras as K
|
||||
import numpy as np
|
||||
import Modules.frequency_spatial_network as fsnet
|
||||
import tensorflow as tf
|
||||
|
||||
under_rate = '20'
|
||||
|
||||
stats = np.load("Data/stats_fs_unet_norm_" + under_rate + ".npy")
|
||||
var_sampling_mask = np.load("Data/sampling_mask_" + under_rate + "perc.npy")
|
||||
|
||||
model = fsnet.wnet(stats[0], stats[1], stats[2], stats[3], kshape = (5,5), kshape2=(3,3))
|
||||
model_name = "Models/wnet_" + under_rate + ".hdf5"
|
||||
model.load_weights(model_name)
|
||||
|
||||
inp = np.random.standard_normal([1, 256, 256, 2]).astype(np.float32)
|
||||
np.save('inp', inp)
|
||||
|
||||
sess = K.backend.get_session()
|
||||
sess.as_default()
|
||||
graph_def = sess.graph.as_graph_def()
|
||||
graph_def = tf.graph_util.convert_variables_to_constants(sess, graph_def, ['conv2d_44/BiasAdd'])
|
||||
with tf.gfile.FastGFile('wnet_20.pb', 'wb') as f:
|
||||
f.write(graph_def.SerializeToString())
|
||||
```
|
||||
|
||||
As a result the TensorFlow frozen model file "wnet_20.pb" is generated.
|
||||
|
||||
### Convert the Frozen TensorFlow Model to Intermediate Representation
|
||||
|
||||
Firstly, open the model in TensorBoard or other TensorFlow model visualization tool. The model supports dynamic
|
||||
batch dimension because the value for the batch dimension is not hardcoded in the model. Model Optimizer need to set all
|
||||
dynamic dimensions to some specific value to create the IR, therefore specify the command line parameter `-b 1` to set
|
||||
the batch dimension equal to 1. The actual batch size dimension can be changed at runtime using the Inference Engine API
|
||||
described in the [Using Shape Inference](../OV_Runtime_UG/ShapeInference.md). Also refer to the General Conversion Parameters section in [Converting a Model to Intermediate Representation (IR)](../MO_DG/prepare_model/convert_model/Converting_Model.md) and [Convert Your TensorFlow* Model](../MO_DG/prepare_model/convert_model/Convert_Model_From_TensorFlow.md)
|
||||
for more details and command line parameters used for the model conversion.
|
||||
|
||||
```sh
|
||||
mo --input_model <PATH_TO_MODEL>/wnet_20.pb -b 1
|
||||
```
|
||||
|
||||
> **NOTE**: This conversion guide is applicable for the 2021.3 release of OpenVINO and starting from 2021.4
|
||||
> OpenVINO has supported this model out of the box.
|
||||
|
||||
Model Optimizer produces the following error:
|
||||
```bash
|
||||
[ ERROR ] List of operations that cannot be converted to Inference Engine IR:
|
||||
[ ERROR ] Complex (1)
|
||||
[ ERROR ] lambda_2/Complex
|
||||
[ ERROR ] IFFT2D (1)
|
||||
[ ERROR ] lambda_2/IFFT2D
|
||||
[ ERROR ] ComplexAbs (1)
|
||||
[ ERROR ] lambda_2/Abs
|
||||
[ ERROR ] Part of the nodes was not converted to IR. Stopped.
|
||||
```
|
||||
|
||||
The error means that the Model Optimizer doesn't know how to handle 3 types of TensorFlow\* operations: "Complex",
|
||||
"IFFT2D" and "ComplexAbs". In order to see more details about the conversion process run the model conversion with
|
||||
additional parameter `--log_level DEBUG`. It is worth to mention the following lines from the detailed output:
|
||||
|
||||
```bash
|
||||
[ INFO ] Called "tf_native_tf_node_infer" for node "lambda_2/Complex"
|
||||
[ <TIMESTAMP> ] [ DEBUG ] [ tf:228 ] Added placeholder with name 'lambda_2/lambda_3/strided_slice_port_0_ie_placeholder'
|
||||
[ <TIMESTAMP> ] [ DEBUG ] [ tf:228 ] Added placeholder with name 'lambda_2/lambda_4/strided_slice_port_0_ie_placeholder'
|
||||
[ <TIMESTAMP> ] [ DEBUG ] [ tf:241 ] update_input_in_pbs: replace input 'lambda_2/lambda_3/strided_slice' with input 'lambda_2/lambda_3/strided_slice_port_0_ie_placeholder'
|
||||
[ <TIMESTAMP> ] [ DEBUG ] [ tf:249 ] Replacing input '0' of the node 'lambda_2/Complex' with placeholder 'lambda_2/lambda_3/strided_slice_port_0_ie_placeholder'
|
||||
[ <TIMESTAMP> ] [ DEBUG ] [ tf:241 ] update_input_in_pbs: replace input 'lambda_2/lambda_4/strided_slice' with input 'lambda_2/lambda_4/strided_slice_port_0_ie_placeholder'
|
||||
[ <TIMESTAMP> ] [ DEBUG ] [ tf:249 ] Replacing input '1' of the node 'lambda_2/Complex' with placeholder 'lambda_2/lambda_4/strided_slice_port_0_ie_placeholder'
|
||||
[ <TIMESTAMP> ] [ DEBUG ] [ tf:148 ] Inferred shape of the output tensor with index '0' of the node 'lambda_2/Complex': '[ 1 256 256]'
|
||||
[ <TIMESTAMP> ] [ DEBUG ] [ infer:145 ] Outputs:
|
||||
[ <TIMESTAMP> ] [ DEBUG ] [ infer:32 ] output[0]: shape = [ 1 256 256], value = <UNKNOWN>
|
||||
[ <TIMESTAMP> ] [ DEBUG ] [ infer:129 ] --------------------
|
||||
[ <TIMESTAMP> ] [ DEBUG ] [ infer:130 ] Partial infer for lambda_2/IFFT2D
|
||||
[ <TIMESTAMP> ] [ DEBUG ] [ infer:131 ] Op: IFFT2D
|
||||
[ <TIMESTAMP> ] [ DEBUG ] [ infer:132 ] Inputs:
|
||||
[ <TIMESTAMP> ] [ DEBUG ] [ infer:32 ] input[0]: shape = [ 1 256 256], value = <UNKNOWN>
|
||||
```
|
||||
|
||||
This is a part of the log of the partial inference phase of the model conversion. See the "Partial Inference" section on
|
||||
the [Model Optimizer Extensibility](../MO_DG/prepare_model/customize_model_optimizer/Customize_Model_Optimizer.md) for
|
||||
more information about this phase. Model Optimizer inferred output shape for the unknown operation of type "Complex"
|
||||
using a "fallback" to TensorFlow. However, it is not enough to generate the IR because Model Optimizer doesn't know
|
||||
which attributes of the operation should be saved to IR. So it is necessary to implement Model Optimizer extensions to
|
||||
support these operations.
|
||||
|
||||
Before going into the extension development it is necessary to understand what these unsupported operations do according
|
||||
to the TensorFlow framework specification.
|
||||
|
||||
* "Complex" - returns a tensor of complex type constructed from two real input tensors specifying real and imaginary
|
||||
part of a complex number.
|
||||
* "IFFT2D" - returns a tensor with inverse 2-dimensional discrete Fourier transform over the inner-most 2 dimensions of
|
||||
an input.
|
||||
* "ComplexAbs" - returns a tensor with absolute values of input tensor with complex numbers.
|
||||
|
||||
The part of the model with all three unsupported operations is depicted below:
|
||||
|
||||
|
||||
|
||||
This model uses complex numbers during the inference but Inference Engine does not support tensors of this data type. So
|
||||
it is necessary to find a way how to avoid using tensors of such a type in the model. Fortunately, the complex tensor
|
||||
appear as a result of "Complex" operation, is used as input in the "IFFT2D" operation then is passed to "ComplexAbs"
|
||||
which produces real value tensor as output. So there are just 3 operations consuming/producing complex tensors in the
|
||||
model.
|
||||
|
||||
Let's design an OpenVINO operation "FFT" which get a single real number tensor describing the complex number and
|
||||
produces a single real number tensor describing output complex tensor. This way the fact that the model uses complex
|
||||
numbers is hidden inside the "FFT" operation implementation. The operation gets a tensor of shape `[N, H, W, 2]` and
|
||||
produces the output tensor with the same shape, where the innermost dimension contains pairs of real numbers describing
|
||||
the complex number (its real and imaginary part). As we will see further this operation will allow us to support the
|
||||
model. The implementation of the Model Optimizer operation should be saved to `mo_extensions/ops/FFT.py` file:
|
||||
|
||||
@snippet FFT.py fft:operation
|
||||
|
||||
The attribute `inverse` is a flag specifying type of the FFT to apply: forward or inverse.
|
||||
|
||||
See the "Model Optimizer Operation" section of [Model Optimizer Extensibility](../MO_DG/prepare_model/customize_model_optimizer/Customize_Model_Optimizer.md) for detailed instructions on how to implement the operation.
|
||||
|
||||
Now it is necessary to implement extractor for the "IFFT2D" operation according to the
|
||||
"Operation Extractor" section of [Model Optimizer Extensibility](../MO_DG/prepare_model/customize_model_optimizer/Customize_Model_Optimizer.md). The
|
||||
following snippet provides two extractors: one for "IFFT2D", another one for "FFT2D", however only on of them is used in this example. The implementation should be saved to the file `mo_extensions/front/tf/FFT_ext.py`.
|
||||
|
||||
@snippet FFT_ext.py fft_ext:extractor
|
||||
|
||||
> **NOTE**: The graph is in inconsistent state after extracting node attributes because according to original operation
|
||||
> "IFFT2D" semantic it should have an input consuming a tensor of complex numbers, but the extractor instantiated an
|
||||
> operation "FFT" which expects a real tensor with specific layout. But the inconsistency will be resolved during
|
||||
> applying front phase transformations discussed below.
|
||||
|
||||
The output shape of the operation "AddV2" from the picture above is `[N, H, W, 2]`. Where the innermost dimension
|
||||
contains pairs of real numbers describing the complex number (its real and imaginary part). The following "StridedSlice"
|
||||
operations split the input tensor into 2 parts to get a tensor of real and a tensor of imaginary parts which are then
|
||||
consumed with the "Complex" operation to produce a tensor of complex numbers. These "StridedSlice" and "Complex"
|
||||
operations can be removed so the "FFT" operation will get a real value tensor encoding complex numbers. To achieve this
|
||||
we implement the front phase transformation which searches for a pattern of two "StridedSlice" operations with specific
|
||||
attributes producing data to "Complex" operation and removes it from the graph. Refer to the
|
||||
"Pattern-Defined Front Phase Transformations" section of [Model Optimizer Extensibility](../MO_DG/prepare_model/customize_model_optimizer/Customize_Model_Optimizer.md) for more
|
||||
information on how this type of transformation works. The code snippet should be saved to the file
|
||||
`mo_extensions/front/tf/Complex.py`.
|
||||
|
||||
@snippet Complex.py complex:transformation
|
||||
|
||||
> **NOTE**: The graph is in inconsistent state because the "ComplexAbs" operation consumes complex value tensor but
|
||||
> "FFT" produces real value tensor.
|
||||
|
||||
Now lets implement a transformation which replace a "ComplexAbs" operation with a sub-graph of primitive operations
|
||||
which calculate the result using the following formulae: \f$module(z) = \sqrt{real(z) \cdot real(z) + imag(z) \cdot imag(z)}\f$.
|
||||
Original "IFFT2D" operation produces tensor of complex values, but the "FFT" operation produces a real value tensor with
|
||||
the same format and shape as the input for the operation. So the input shape for the "ComplexAbs" will be `[N, H, W, 2]`
|
||||
with the innermost dimension containing tuple with real and imaginary part of a complex number. In order to calculate
|
||||
absolute values for the complex tensor we do the following:
|
||||
1. Raise all elements in the power of 2.
|
||||
2. Calculate a reduced sum over the innermost dimension.
|
||||
3. Calculate a square root.
|
||||
|
||||
The implementation should be saved to the file `mo_extensions/front/tf/ComplexAbs.py` and provided below:
|
||||
|
||||
@snippet ComplexAbs.py complex_abs:transformation
|
||||
|
||||
Now it is possible to convert the model using the following command line:
|
||||
```sh
|
||||
mo --input_model <PATH_TO_MODEL>/wnet_20.pb -b 1 --extensions mo_extensions/
|
||||
```
|
||||
|
||||
The sub-graph corresponding to the originally non-supported one is depicted in the image below:
|
||||
|
||||
|
||||
|
||||
> **NOTE**: Model Optimizer performed conversion of the model from NHWC to NCHW layout that is why the dimension with
|
||||
> the value 2 moved to another position.
|
||||
|
||||
### Inference Engine Extension Implementation
|
||||
Now it is necessary to implement the extension for the CPU plugin with operation "FFT" introduced previously. The code
|
||||
below is based on the template extension described in [Inference Engine Extensibility Mechanism](../OV_Runtime_UG/Extensibility_DG/Intro.md).
|
||||
|
||||
#### CMake Build File
|
||||
The first step is to create a CMake configuration file which builds the extension. The content of the "CMakeLists.txt"
|
||||
file is the following:
|
||||
|
||||
@snippet template_extension/old/CMakeLists.txt cmake:extension
|
||||
|
||||
The CPU FFT kernel implementation uses OpenCV to perform the FFT that is why the extension library is linked with
|
||||
`opencv_core` which comes with the OpenVINO.
|
||||
|
||||
#### Custom nGraph Operation "FFT" Implementation
|
||||
The next step is to create the nGraph operation FFT. The header file "fft_op.hpp" has the following content:
|
||||
|
||||
@snippet template_extension/old/fft_op.hpp fft_op:header
|
||||
|
||||
The operation has just one boolean attribute `inverse`. Implementation of the necessary nGraph operation functions are
|
||||
in the `fft_op.cpp` file with the following content:
|
||||
|
||||
@snippet template_extension/old/fft_op.cpp fft_op:implementation
|
||||
|
||||
Refer to the [Custom nGraph Operation](../OV_Runtime_UG/Extensibility_DG/AddingNGraphOps.md) for more details.
|
||||
|
||||
#### CPU FFT Kernel Implementation
|
||||
The operation implementation for CPU plugin uses OpenCV to perform the FFT. The header file "fft_kernel.hpp" has the
|
||||
following content:
|
||||
|
||||
@snippet template_extension/old/fft_kernel.hpp fft_kernel:header
|
||||
|
||||
The "fft_kernel.cpp" with the implementation of the CPU has the following content:
|
||||
|
||||
@snippet template_extension/old/fft_kernel.cpp fft_kernel:implementation
|
||||
|
||||
Refer to the [How to Implement Custom CPU Operations](../OV_Runtime_UG/Extensibility_DG/CPU_Kernel.md) for more details.
|
||||
|
||||
#### Extension Library Implementation
|
||||
The last step is to create an extension library "extension.cpp" and "extension.hpp" which will include the FFT
|
||||
operation for the CPU plugin. The code of the library is described in the [Extension Library](../OV_Runtime_UG/Extensibility_DG/Extension.md).
|
||||
|
||||
### Building and Running the Custom Extension
|
||||
To build the extension, run the following:<br>
|
||||
```bash
|
||||
mkdir build && cd build
|
||||
source /opt/intel/openvino_2022/setupvars.sh
|
||||
cmake .. -DCMAKE_BUILD_TYPE=Release
|
||||
make --jobs=$(nproc)
|
||||
```
|
||||
|
||||
The result of this command is a compiled shared library (`.so` or `.dll`). It should be loaded in the
|
||||
application using `Core` class instance method `AddExtension` like this
|
||||
`core.AddExtension(std::make_shared<Extension>(compiled_library_file_name), "CPU");`.
|
||||
|
||||
To test that the extension is implemented correctly we can run the "mri_reconstruction_demo" with the following content:
|
||||
|
||||
@snippet mri_reconstruction_demo.py mri_demo:demo
|
||||
|
||||
The script can be executed using the following command line:
|
||||
```bash
|
||||
python3 mri_reconstruction_demo.py \
|
||||
-m <PATH_TO_IR>/wnet_20.xml \
|
||||
-i <PATH_TO_SAMPLE_MRI_IMAGE>.npy \
|
||||
-p <Hybrid-CS-Model-MRI_repo>/Data/sampling_mask_20perc.npy \
|
||||
-l <PATH_TO_BUILD_DIR>/libtemplate_extension.so \
|
||||
-d CPU
|
||||
```
|
||||
|
||||
## Additional Resources
|
||||
|
||||
- Intel® Distribution of OpenVINO™ toolkit home page: [https://software.intel.com/en-us/openvino-toolkit](https://software.intel.com/en-us/openvino-toolkit)
|
||||
- OpenVINO™ toolkit online documentation: [https://docs.openvino.ai](https://docs.openvino.ai)
|
||||
- [Model Optimizer Developer Guide](../MO_DG/Deep_Learning_Model_Optimizer_DevGuide.md)
|
||||
- [Model Optimizer Extensibility](../MO_DG/prepare_model/customize_model_optimizer/Customize_Model_Optimizer.md)
|
||||
- [Inference Engine Extensibility Mechanism](../OV_Runtime_UG/Extensibility_DG/Intro.md)
|
||||
- [OpenVINO™ Toolkit Samples Overview](../OV_Runtime_UG/Samples_Overview.md)
|
||||
- [Overview of OpenVINO™ Toolkit Pre-Trained Models](@ref omz_models_group_intel)
|
||||
- For IoT Libraries and Code Samples see the [Intel® IoT Developer Kit](https://github.com/intel-iot-devkit).
|
||||
|
||||
## Converting Models:
|
||||
|
||||
- [Convert Your Caffe Model](../MO_DG/prepare_model/convert_model/Convert_Model_From_Caffe.md)
|
||||
- [Convert Your TensorFlow Model](../MO_DG/prepare_model/convert_model/Convert_Model_From_TensorFlow.md)
|
||||
- [Convert Your MXNet Model](../MO_DG/prepare_model/convert_model/Convert_Model_From_MxNet.md)
|
||||
- [Convert Your Kaldi Model](../MO_DG/prepare_model/convert_model/Convert_Model_From_Kaldi.md)
|
||||
- [Convert Your ONNX Model](../MO_DG/prepare_model/convert_model/Convert_Model_From_ONNX.md)
|
||||
- [Convert Your PaddlePaddle Model](../MO_DG/prepare_model/convert_model/Convert_Model_From_Paddle.md)
|
||||
@@ -214,7 +214,7 @@ One of the layers in the specified topology might not have inputs or values. Ple
|
||||
|
||||
#### 24. What does the message "Part of the nodes was not translated to IE. Stopped" mean? <a name="question-24"></a>
|
||||
|
||||
Some of the layers are not supported by the Inference Engine and cannot be translated to an Intermediate Representation. You can extend the Model Optimizer by allowing generation of new types of layers and implement these layers in the dedicated Inference Engine plugins. For more information, refer to the [Custom Layers Guide](../../HOWTO/Custom_Layers_Guide.md) and [Inference Engine Extensibility Mechanism](../../OV_Runtime_UG/Extensibility_DG/Intro.md)
|
||||
Some of the layers are not supported by the Inference Engine and cannot be translated to an Intermediate Representation. You can extend the Model Optimizer by allowing generation of new types of layers and implement these layers in the dedicated Inference Engine plugins. For more information, refer to the [OpenVINO™ Extensibility Mechanism](../../Extensibility_UG/Intro.md)
|
||||
|
||||
#### 25. What does the message "While creating an edge from .. to .. : node name is undefined in the graph. Check correctness of the input model" mean? <a name="question-25"></a>
|
||||
|
||||
@@ -268,7 +268,7 @@ Model Optimizer tried to write an event file in the specified directory but fail
|
||||
|
||||
#### 37. What does the message "There is no registered 'infer' function for node with op = .. . Please implement this function in the extensions" mean? <a name="question-37"></a>
|
||||
|
||||
Most likely, you tried to extend Model Optimizer with a new primitive, but did not specify an infer function. For more information on extensions, see [Custom Layers Guide](../../HOWTO/Custom_Layers_Guide.md).
|
||||
Most likely, you tried to extend Model Optimizer with a new primitive, but did not specify an infer function. For more information on extensions, see [OpenVINO™ Extensibility Mechanism](../../Extensibility_UG/Intro.md).
|
||||
|
||||
#### 38. What does the message "Stopped shape/value propagation at node" mean? <a name="question-38"></a>
|
||||
|
||||
@@ -300,7 +300,7 @@ Most likely, there is a problem with the specified file for model. The file exis
|
||||
|
||||
#### 45. What does the message "Found custom layer. Model Optimizer does not support this layer. Please, register it in CustomLayersMapping.xml or implement extension" mean? <a name="question-45"></a>
|
||||
|
||||
This means that the layer `{layer_name}` is not supported in the Model Optimizer. You can find a list of all unsupported layers in the corresponding section. You should implement the extensions for this layer ([Custom Layers Guide](../../HOWTO/Custom_Layers_Guide.md)).
|
||||
This means that the layer `{layer_name}` is not supported in the Model Optimizer. You can find a list of all unsupported layers in the corresponding section. You should implement the extensions for this layer ([OpenVINO™ Extensibility Mechanism](../../Extensibility_UG/Intro.md)).
|
||||
|
||||
#### 46. What does the message "Custom replacement configuration file does not exist" mean? <a name="question-46"></a>
|
||||
|
||||
@@ -308,7 +308,7 @@ Path to the custom replacement configuration file was provided with the `--trans
|
||||
|
||||
#### 47. What does the message "Extractors collection have case insensitive duplicates" mean? <a name="question-47"></a>
|
||||
|
||||
When extending Model Optimizer with new primitives keep in mind that their names are case insensitive. Most likely, another operation with the same name is already defined. For more information, see [Custom Layers Guide](../../HOWTO/Custom_Layers_Guide.md).
|
||||
When extending Model Optimizer with new primitives keep in mind that their names are case insensitive. Most likely, another operation with the same name is already defined. For more information, see [OpenVINO™ Extensibility Mechanism](../../Extensibility_UG/Intro.md).
|
||||
|
||||
#### 48. What does the message "Input model name is not in an expected format, cannot extract iteration number" mean? <a name="question-48"></a>
|
||||
|
||||
@@ -340,7 +340,7 @@ Please, make sure that inputs are defined and have correct shapes. You can use `
|
||||
|
||||
#### 55. What does the message "Attempt to register of custom name for the second time as class. Note that custom names are case-insensitive" mean? <a name="question-55"></a>
|
||||
|
||||
When extending Model Optimizer with new primitives keep in mind that their names are case insensitive. Most likely, another operation with the same name is already defined. For more information, see [Custom Layers Guide](../../HOWTO/Custom_Layers_Guide.md).
|
||||
When extending Model Optimizer with new primitives keep in mind that their names are case insensitive. Most likely, another operation with the same name is already defined. For more information, see [OpenVINO™ Extensibility Mechanism](../../Extensibility_UG/Intro.md).
|
||||
|
||||
#### 56. What does the message "Both --input_shape and --batch were provided. Please, provide only one of them" mean? <a name="question-56"></a>
|
||||
|
||||
@@ -492,7 +492,7 @@ For more information, refer to [Converting a MXNet* Model](convert_model/Convert
|
||||
|
||||
Model Optimizer tried to load the model that contains some unsupported operations.
|
||||
If you want to convert model that contains unsupported operations you need to prepare extension for all such operations.
|
||||
For more information, refer to [Custom Layers Guide](../../HOWTO/Custom_Layers_Guide.md).
|
||||
For more information, refer to [OpenVINO™ Extensibility Mechanism](../../Extensibility_UG/Intro.md).
|
||||
|
||||
#### 87. What does the message "Can not register Op ... Please, call function 'register_caffe_python_extractor' with parameter 'name'" mean? <a name="question-87"></a>
|
||||
|
||||
@@ -538,7 +538,7 @@ Note that the first call <code>register_caffe_python_extractor(ProposalPythonExa
|
||||
|
||||
The second call prevents Model Optimizer from using this extension as if it is an extension for
|
||||
a layer with type `Proposal`. Otherwise, this layer can be chosen as an implementation of extension that can lead to potential issues.
|
||||
For more information, refer to the [Custom Layers Guide](../../HOWTO/Custom_Layers_Guide.md).
|
||||
For more information, refer to the [OpenVINO™ Extensibility Mechanism](../../Extensibility_UG/Intro.md).
|
||||
|
||||
#### 88. What does the message "Model Optimizer is unable to calculate output shape of Memory node .." mean? <a name="question-88"></a>
|
||||
|
||||
@@ -573,7 +573,7 @@ This message means that if you have model with custom layers and its json file h
|
||||
lower than 1.0.0, Model Optimizer does not support such topologies. If you want to convert it you have to rebuild
|
||||
MXNet with unsupported layers or generate new json with MXNet version 1.0.0 and higher. Also you need to implement
|
||||
Inference Engine extension for used custom layers.
|
||||
For more information, refer to the [Custom Layers Guide](../../HOWTO/Custom_Layers_Guide.md).
|
||||
For more information, refer to the [OpenVINO™ Extensibility Mechanism](../../Extensibility_UG/Intro.md).
|
||||
|
||||
#### 97. What does the message "Graph contains a cycle. Can not proceed .." mean? <a name="question-97"></a>
|
||||
|
||||
@@ -586,7 +586,7 @@ For Tensorflow:
|
||||
|
||||
For all frameworks:
|
||||
1. [Replace cycle containing Sub-graph in Model Optimizer](customize_model_optimizer/Subgraph_Replacement_Model_Optimizer.md)
|
||||
2. [Custom Layers Guide](../../HOWTO/Custom_Layers_Guide.md)
|
||||
2. [OpenVINO™ Extensibility Mechanism](../../Extensibility_UG/Intro.md)
|
||||
|
||||
or
|
||||
* Edit network in original framework to exclude cycle.
|
||||
|
||||
@@ -354,7 +354,7 @@ TensorFlow*-specific parameters:
|
||||
mo --input_model inception_v1.pb -b 1 --tensorboard_logdir /tmp/log_dir --output_dir <OUTPUT_MODEL_DIR>
|
||||
```
|
||||
|
||||
* Launching the Model Optimizer for a model with custom TensorFlow operations (refer to the [TensorFlow* documentation](https://www.tensorflow.org/extend/adding_an_op)) implemented in C++ and compiled into the shared library `my_custom_op.so`. Model Optimizer falls back to TensorFlow to infer output shape of operations implemented in the library if a custom TensorFlow operation library is provided. If it is not provided, a custom operation with an inference function is needed. For more information about custom operations, refer to the [Custom Layers Guide](../../../HOWTO/Custom_Layers_Guide.md).
|
||||
* Launching the Model Optimizer for a model with custom TensorFlow operations (refer to the [TensorFlow* documentation](https://www.tensorflow.org/extend/adding_an_op)) implemented in C++ and compiled into the shared library `my_custom_op.so`. Model Optimizer falls back to TensorFlow to infer output shape of operations implemented in the library if a custom TensorFlow operation library is provided. If it is not provided, a custom operation with an inference function is needed. For more information about custom operations, refer to the [OpenVINO™ Extensibility Mechanism](../../../Extensibility_UG/Intro.md).
|
||||
```sh
|
||||
mo --input_model custom_model.pb --tensorflow_custom_layer_libraries ./my_custom_op.so --output_dir <OUTPUT_MODEL_DIR>
|
||||
```
|
||||
|
||||
@@ -269,7 +269,7 @@ mo --input_model bvlc_alexnet.caffemodel --reverse_input_channels --mean_values
|
||||
```
|
||||
|
||||
Launch the Model Optimizer for the Caffe bvlc_alexnet model with extensions listed in specified directories, specified mean_images binaryproto
|
||||
file. For more information about extensions, please refer to the [Custom Layers Guide](../../../HOWTO/Custom_Layers_Guide.md).
|
||||
file. For more information about extensions, please refer to the [OpenVINO™ Extensibility Mechanism](../../../Extensibility_UG/Intro.md).
|
||||
```sh
|
||||
mo --input_model bvlc_alexnet.caffemodel --extensions /home/,/some/other/path/ --mean_file /path/to/binaryproto --output_dir <OUTPUT_MODEL_DIR>
|
||||
```
|
||||
|
||||
@@ -1261,5 +1261,5 @@ Refer to the `extensions/back/GatherNormalizer.py` for the example of a such typ
|
||||
* [Deep Learning Network Intermediate Representation and Operation Sets in OpenVINO™](../../IR_and_opsets.md)
|
||||
* [Converting a Model to Intermediate Representation (IR)](../convert_model/Converting_Model.md)
|
||||
* [OpenVINO Model Representation](../../../OV_Runtime_UG/model_representation.md)
|
||||
* [Inference Engine Extensibility Mechanism](../../../OV_Runtime_UG/Extensibility_DG/Intro.md)
|
||||
* [OpenVINO™ Extensibility Mechanism](../../../Extensibility_UG/Intro.md)
|
||||
* [Extending the Model Optimizer with Caffe* Python Layers](Extending_Model_Optimizer_with_Caffe_Python_Layers.md)
|
||||
@@ -1,82 +0,0 @@
|
||||
# Custom nGraph Operations {#openvino_docs_IE_DG_Extensibility_DG_AddingNGraphOps}
|
||||
|
||||
Inference Engine Extension API allows you to register operation sets (opsets) with custom nGraph operations to support models with operations which OpenVINO™ does not support out-of-the-box.
|
||||
|
||||
Besides creating custom nGraph operations, to [support custom operations](../../HOWTO/Custom_Layers_Guide.md) in your model you must also create a Model Optimizer extension for the custom operations and an Inference Engine device plugin extension for the device you will use for inference.
|
||||
|
||||
## Operation Class
|
||||
|
||||
To add your custom nGraph operation, create a new class that extends `ngraph::Op`, which is in turn derived from `ngraph::Node`, the base class for all graph operations in nGraph. Follow the steps below to add a custom nGraph operation:
|
||||
|
||||
1. Add the `NGRAPH_RTTI_DECLARATION` and `NGRAPH_RTTI_DEFINITION` macros which define a `NodeTypeInfo` object that identifies the type of the operation to the graph users and helps with dynamic type resolution. The type info of an nGraph operation currently consists of a string identifier and a version number, but this may change in the future.
|
||||
|
||||
2. Implement constructors that optionally take the operation inputs and attributes as parameters.
|
||||
|
||||
3. Override the shape inference method `validate_and_infer_types`. This method is called multiple times during graph manipulations to determine the shapes and element types of the operations outputs. To access the input shapes and input element types, use the `get_input_partial_shape()` and `get_input_element_type()` methods of `ngraph::Node`. Set the inferred shape and element type of the output using `set_output_type`.
|
||||
|
||||
4. Override the `clone_with_new_inputs` method, which enables graph manipulation routines to create copies of this operation and connect it to different nodes during optimization.
|
||||
|
||||
5. Override the `visit_attributes` method, which enables serialization and deserialization of operation attributes. An `AttributeVisitor` is passed to the method, and the implementation is expected to walk over all the attributes in the op using the type-aware `on_attribute` helper. Helpers are already implemented for standard C++ types like `int64_t`, `float`, `bool`, `vector`, and for existing nGraph defined types.
|
||||
|
||||
6. Override `evaluate`, which is an optional method that enables the application of constant folding if there is a custom operation on the constant branch. If your operation contains `evaluate` method you also need to override the `has_evaluate` method, this method allow to get information about availability of `evaluate` method for the operation.
|
||||
|
||||
Based on that, declaration of an operation class can look as follows:
|
||||
|
||||
@snippet template_extension/old/op.hpp op:header
|
||||
|
||||
### Class Fields
|
||||
|
||||
The provided implementation has several fields:
|
||||
|
||||
* `add` of type `int64_t` is an attribute of a custom operation
|
||||
* `type_info` of type `ngraph::NodeTypeInfo` defines type and version of an operation
|
||||
|
||||
### Operation Constructors
|
||||
|
||||
nGraph operation contains two constructors:
|
||||
* Default constructor, which enables you to create an operation without attributes
|
||||
* Constructor that creates and validates an operation with specified inputs and attributes
|
||||
|
||||
@snippet template_extension/old/op.cpp op:ctor
|
||||
|
||||
### `validate_and_infer_types()`
|
||||
|
||||
`ngraph::Node::validate_and_infer_types` method validates operation attributes and calculates output shapes using attributes of the operation.
|
||||
|
||||
@snippet template_extension/old/op.cpp op:validate
|
||||
|
||||
### `clone_with_new_inputs()`
|
||||
|
||||
`ngraph::Node::clone_with_new_inputs` method creates a copy of the nGraph operation with new inputs.
|
||||
|
||||
@snippet template_extension/old/op.cpp op:copy
|
||||
|
||||
### `visit_attributes()`
|
||||
|
||||
`ngraph::Node::visit_attributes` method enables you to visit all operation attributes.
|
||||
|
||||
@snippet template_extension/old/op.cpp op:visit_attributes
|
||||
|
||||
### `evaluate()` and `has_evaluate()`
|
||||
|
||||
`ngraph::Node::evaluate` method enables you to apply constant folding to an operation.
|
||||
|
||||
@snippet template_extension/old/op.cpp op:evaluate
|
||||
|
||||
## Register Custom Operations in Extension Class
|
||||
|
||||
To add custom operations to the [Extension](Extension.md) class, create an operation set with custom operations and implement the `InferenceEngine::IExtension::getOpSets` method:
|
||||
|
||||
@snippet template_extension/old/extension.cpp extension:getOpSets
|
||||
|
||||
This method returns a map of opsets that exist in the [extension library](Extension.md).
|
||||
nGraph provides an opset mechanism to group operations into clusters. Different opsets distinguish between different versions of one operation.
|
||||
|
||||
When specifying opset names, follow the rules below:
|
||||
* Use unique opset names.
|
||||
* Do not use the following built-in opset names: `extension`, `experimental`, `opset1`, `opset2`, `opset3`, ... , `opsetN`.
|
||||
* [Make sure that the Model Optimizer](../../HOWTO/Custom_Layers_Guide.md) and your extension use the same opset names.
|
||||
* IR v10 operations have the mandatory `version` attribute specifying the opset.
|
||||
Operations from the default opset cannot be redefined.
|
||||
|
||||
Use a custom opset to create a new operation or extend functionality of an existing operation from another opset.
|
||||
@@ -1,19 +0,0 @@
|
||||
# Build Extension Library Using CMake* {#openvino_docs_IE_DG_Extensibility_DG_Building}
|
||||
|
||||
Inference Engine build infrastructure provides the Inference Engine Package for application development.
|
||||
|
||||
To configure the build of your extension library, use the following CMake script:
|
||||
|
||||
@snippet template_extension/old/CMakeLists.txt cmake:extension
|
||||
|
||||
This CMake script finds the Inference Engine and nGraph using the `find_package` CMake command.
|
||||
|
||||
To build the extension library, run the commands below:
|
||||
|
||||
```sh
|
||||
$ cd template_extension/old
|
||||
$ mkdir build
|
||||
$ cd build
|
||||
$ cmake -DOpenVINO_DIR=[OpenVINO_DIR] ../
|
||||
$ cmake --build .
|
||||
```
|
||||
@@ -1,71 +0,0 @@
|
||||
# CPU Kernel Custom Operations {#openvino_docs_IE_DG_Extensibility_DG_CPU_Kernel}
|
||||
|
||||
To enable operations not supported by OpenVINO™ out of the box, you need a custom extension for Model Optimizer, a custom nGraph operation set, and a custom kernel for the device you will target. This page describes custom kernel support for the CPU device.
|
||||
|
||||
The primary means of the performance of the CPU codepath in the Inference Engine is the Intel® Math Kernel Library for Deep Neural Networks (Intel® MKL-DNN), and new CPU kernels extend the Inference Engine plugin for the Intel MKL-DNN. Implementing the InferenceEngine::ILayerExecImpl API call defines a general CPU-side extension. There are no Intel MKL-DNN specifics in the way you need to implement a kernel.
|
||||
|
||||
## Implementation Class
|
||||
|
||||
All custom kernels for the CPU plugin should be inherited from the InferenceEngine::ILayerExecImpl interface.
|
||||
Based on that, declaration of a kernel implementation class can look as follows:
|
||||
|
||||
@snippet template_extension/old/cpu_kernel.hpp cpu_implementation:header
|
||||
|
||||
### Class Fields
|
||||
|
||||
The provided implementation has several fields:
|
||||
|
||||
* `add` of the type `int64_t` is an attribute of a custom operation.
|
||||
* `inShape` of the type `ngraph::Shape` is an input shape.
|
||||
* `outShape` of the type `ngraph::Shape` is an output shape.
|
||||
* `error` of the type `std::string` is a field to handle errors from a constructor.
|
||||
|
||||
### Constructor of Implementation
|
||||
|
||||
An implementation constructor checks parameters of an nGraph operation, stores required attributes, and stores an error message in case of an error.
|
||||
|
||||
@snippet template_extension/old/cpu_kernel.cpp cpu_implementation:ctor
|
||||
|
||||
### `getSupportedConfigurations`
|
||||
|
||||
The InferenceEngine::ILayerExecImpl::getSupportedConfigurations method returns all supported configuration formats (input/output tensor layouts) for your implementation. To specify formats of data, use InferenceEngine::TensorDesc.
|
||||
|
||||
@snippet template_extension/old/cpu_kernel.cpp cpu_implementation:getSupportedConfigurations
|
||||
|
||||
### `init`
|
||||
|
||||
The InferenceEngine::ILayerExecImpl::init method gets a runtime-selected configuration from a vector that is populated from the `getSupportedConfigurations` method and checks the parameters:
|
||||
|
||||
@snippet template_extension/old/cpu_kernel.cpp cpu_implementation:init
|
||||
|
||||
### `execute`
|
||||
|
||||
The InferenceEngine::ILayerExecImpl::execute method accepts and processes the actual tensors as input/output blobs:
|
||||
|
||||
@snippet template_extension/old/cpu_kernel.cpp cpu_implementation:execute
|
||||
|
||||
## Register Implementation in `Extension` Class
|
||||
|
||||
To register custom kernel implementation in the [Extension](Extension.md) class, implement the following methods:
|
||||
|
||||
* <a href="#getImpTypes">getImplTypes</a>
|
||||
* <a href="#getImplementation">getImplementation</a>
|
||||
|
||||
### <a name="getImpTypes"><code>getImplTypes</code></a>
|
||||
|
||||
InferenceEngine::IExtension::getImplTypes returns a vector of implementation types for an operation.
|
||||
|
||||
@snippet template_extension/old/extension.cpp extension:getImplTypes
|
||||
|
||||
### <a name="getImplementation"><code>getImplementation</code></a>
|
||||
|
||||
InferenceEngine::IExtension::getImplementation returns the kernel implementation with a specified type for an operation.
|
||||
|
||||
@snippet template_extension/old/extension.cpp extension:getImplementation
|
||||
|
||||
|
||||
## Load Extension with Executable Kernels to Plugin
|
||||
|
||||
Use the `AddExtension` method of the general plugin interface to load your primitives:
|
||||
|
||||
@snippet snippets/CPU_Kernel.cpp part0
|
||||
@@ -1,78 +0,0 @@
|
||||
# Custom ONNX* Operators {#openvino_docs_IE_DG_Extensibility_DG_Custom_ONNX_Ops}
|
||||
|
||||
The ONNX\* importer provides a mechanism to register custom ONNX operators based on predefined or custom nGraph operations.
|
||||
The function responsible for registering a new operator is called `ngraph::onnx_import::register_operator` and defined in the `onnx_import/onnx_utils.hpp` file.
|
||||
|
||||
## Register Custom ONNX Operator Based on Predefined nGraph Operations
|
||||
|
||||
The steps below explain how to register a custom ONNX operator, for example, CustomRelu, in a domain called `com.example`.
|
||||
CustomRelu is defined as follows:
|
||||
```
|
||||
x >= 0 => f(x) = x * alpha
|
||||
x < 0 => f(x) = x * beta
|
||||
```
|
||||
where `alpha` and `beta` are float constants.
|
||||
|
||||
1. Include headers:
|
||||
|
||||
@snippet onnx_custom_op/onnx_custom_op.cpp onnx_custom_op:headers
|
||||
|
||||
2. Register the CustomRelu operator in the ONNX importer:
|
||||
|
||||
@snippet onnx_custom_op/onnx_custom_op.cpp onnx_custom_op:register_operator
|
||||
|
||||
The `register_operator` function takes four arguments: op_type, opset version, domain, and a function object.
|
||||
The function object is a user-defined function that takes `ngraph::onnx_import::Node` as an input and based on that, returns a graph with nGraph operations.
|
||||
The `ngraph::onnx_import::Node` class represents a node in an ONNX model. It provides functions to fetch input node(s) using `get_ng_inputs`, attribute value using `get_attribute_value`, and many more. See the `onnx_import/core/node.hpp` file for the full class declaration.
|
||||
|
||||
New operator registration must happen before an ONNX model is read. For example, if an model uses the `CustomRelu` operator, call `register_operator("CustomRelu", ...)` before InferenceEngine::Core::ReadNetwork.
|
||||
Reregistering ONNX operators within the same process is supported. If you register an existing operator, you get a warning.
|
||||
|
||||
The example below demonstrates an exemplary model that requires a previously created `CustomRelu` operator:
|
||||
```
|
||||
@include onnx_custom_op/custom_relu_model.prototxt
|
||||
```
|
||||
|
||||
This model is in text format, so before it can be passed to Inference Engine, it has to be converted to binary using:
|
||||
```py
|
||||
from google.protobuf import text_format
|
||||
import onnx
|
||||
|
||||
with open("custom_relu_model.prototxt") as in_file:
|
||||
proto = onnx.ModelProto()
|
||||
text_format.Parse(in_file.read(), proto, allow_field_number=True)
|
||||
s = onnx._serialize(proto)
|
||||
onnx._save_bytes(s, "custom_relu_model.onnx")
|
||||
```
|
||||
|
||||
|
||||
To create a graph with nGraph operations, visit [Custom nGraph Operations](AddingNGraphOps.md).
|
||||
For a complete list of predefined nGraph operators, visit [Available Operations Sets](../../ops/opset.md).
|
||||
|
||||
If you do not need an operator anymore, unregister it by calling `unregister_operator`. The function takes three arguments: `op_type`, `version`, and `domain`.
|
||||
|
||||
@snippet onnx_custom_op/onnx_custom_op.cpp onnx_custom_op:unregister_operator
|
||||
|
||||
## Register Custom ONNX Operator Based on Custom nGraph Operations
|
||||
|
||||
The same principles apply when registering a custom ONNX operator based on custom nGraph operations.
|
||||
This example shows how to register a custom ONNX operator based on `Operation` presented in [this tutorial](AddingNGraphOps.md), which is used in [TemplateExtension](Extension.md):
|
||||
|
||||
@snippet template_extension/old/extension.cpp extension:ctor
|
||||
|
||||
Here, the `register_operator` function is called in the constructor of Extension. The constructor makes sure that the function is called before InferenceEngine::Core::ReadNetwork, because InferenceEngine::Core::AddExtension must be called before a model with a custom operator is read.
|
||||
|
||||
The example below demonstrates how to unregister an operator from the destructor of Extension:
|
||||
@snippet template_extension/old/extension.cpp extension:dtor
|
||||
|
||||
> **REQUIRED**: It is mandatory to unregister a custom ONNX operator if it is defined in a dynamic shared library.
|
||||
|
||||
## Requirements for Building with CMake
|
||||
|
||||
A program that uses the `register_operator` functionality requires `openvino::core` and `openvino::frontend::onnx` libraries in addition to the OpenVINO Inference Runtime.
|
||||
The `openvino::frontend::onnx` is a component of the `OpenVINO` package , so `find_package(OpenVINO REQUIRED COMPONENTS ONNX)` can find both.
|
||||
Those libraries need to be passed to the `target_link_libraries` command in the CMakeLists.txt file.
|
||||
|
||||
See CMakeLists.txt below for reference:
|
||||
|
||||
@snippet onnx_custom_op/CMakeLists.txt cmake:onnx_custom_op
|
||||
@@ -1,29 +0,0 @@
|
||||
# Extension Library {#openvino_docs_IE_DG_Extensibility_DG_Extension}
|
||||
|
||||
Inference Engine provides an InferenceEngine::IExtension interface, which defines the interface for Inference Engine Extension libraries.
|
||||
Inherit all extension libraries from this interface. The example below contains an implementation of two operations: `Template`
|
||||
used as an example in this document and `FFT` used as a more complex example from the [Custom Operations Guide](../../HOWTO/Custom_Layers_Guide.md).
|
||||
|
||||
> **NOTE**: `FFT` operation is implemented using the OpenCV library functions `cv::dft` and `cv::idft`.
|
||||
|
||||
Based on that, the declaration of an extension class can look as follows:
|
||||
|
||||
@snippet template_extension/old/extension.hpp extension:header
|
||||
|
||||
The extension library should use `IE_DEFINE_EXTENSION_CREATE_FUNCTION` macro to export a function, which creates an `Extension` class:
|
||||
|
||||
@snippet template_extension/old/extension.cpp extension:CreateExtension
|
||||
|
||||
Also, an `Extension` object should implement the following methods:
|
||||
|
||||
* InferenceEngine::IExtension::Release deletes an extension object.
|
||||
|
||||
* InferenceEngine::IExtension::GetVersion returns information about the version of the library.
|
||||
|
||||
@snippet template_extension/old/extension.cpp extension:GetVersion
|
||||
|
||||
Implement the InferenceEngine::IExtension::getOpSets method if the extension contains custom layers.
|
||||
Read [Custom nGraph Operation](AddingNGraphOps.md) for more information.
|
||||
|
||||
To integrate execution kernels to the extension library, read [How to Implement Custom CPU Operations](CPU_Kernel.md).
|
||||
To register a custom ONNX\* operator to the extension library, read [Custom ONNX Operators](Custom_ONNX_Ops.md).
|
||||
@@ -1,233 +0,0 @@
|
||||
# How to Implement Custom GPU Operations {#openvino_docs_IE_DG_Extensibility_DG_GPU_Kernel}
|
||||
|
||||
To enable operations not supported by OpenVINO™ out of the box, you need a custom extension for Model Optimizer, a custom nGraph operation set, and a custom kernel for the device you will target. This page describes custom kernel support for the GPU device.
|
||||
|
||||
The GPU codepath abstracts many details about OpenCL\*. You need to provide the kernel code in OpenCL C and an XML configuration file that connects the kernel and its parameters to the parameters of the operation.
|
||||
|
||||
There are two options for using the custom operation configuration file:
|
||||
|
||||
* Include a section with your kernels into the global automatically-loaded `cldnn_global_custom_kernels/cldnn_global_custom_kernels.xml` file, which is hosted in the `<INSTALL_DIR>/runtime/bin` folder
|
||||
* Call the `InferenceEngine::Core::SetConfig()` method from your application with the `InferenceEngine::PluginConfigParams::KEY_CONFIG_FILE` key and the configuration file name as a value before loading the network that uses custom operations to the plugin:
|
||||
|
||||
@snippet snippets/GPU_Kernel.cpp part0
|
||||
|
||||
All Inference Engine samples, except the trivial `hello_classification`, and most Open Model Zoo demos
|
||||
feature a dedicated command-line option `-c` to load custom kernels. For example, to load custom operations for the classification sample, run the command below:
|
||||
```sh
|
||||
$ ./classification_sample -m <path_to_model>/bvlc_alexnet_fp16.xml -i ./validation_set/daily/227x227/apron.bmp -d GPU
|
||||
-c <absolute_path_to_config>/custom_layer_example.xml
|
||||
```
|
||||
|
||||
## Configuration File Format <a name="config-file-format"></a>
|
||||
|
||||
The configuration file is expected to follow the `.xml` file structure
|
||||
with a node of the type `CustomLayer` for every custom operation you provide.
|
||||
|
||||
The definitions described in the sections below use the following notations:
|
||||
|
||||
Notation | Description
|
||||
---|---
|
||||
(0/1) | Can have zero or one instance of this node or attribute
|
||||
(1) | Must have only one instance of this node or attribute
|
||||
(0+) | Can have any number of instances of this node or attribute
|
||||
(1+) | Can have one or more instances of this node or attribute
|
||||
|
||||
### CustomLayer Node and Sub-Node Structure
|
||||
|
||||
`CustomLayer` node contains the entire configuration for a single custom operation.
|
||||
|
||||
| Attribute Name |\# | Description |
|
||||
|-----|-----|-----|
|
||||
| `name` | (1) | The name of the operation type to be used. This name should be identical to the type used in the IR.|
|
||||
| `type` | (1) | Must be `SimpleGPU`. |
|
||||
| `version` | (1) | Must be `1`. |
|
||||
|
||||
**Sub-nodes**: `Kernel` (1), `Buffers` (1), `CompilerOptions` (0+),
|
||||
`WorkSizes` (0/1)
|
||||
|
||||
### Kernel Node and Sub-Node Structure
|
||||
|
||||
`Kernel` node contains all kernel source code configuration. No kernel
|
||||
node structure exists.
|
||||
|
||||
**Sub-nodes**: `Source` (1+), `Define` (0+)
|
||||
|
||||
### Source Node and Sub-Node Structure
|
||||
|
||||
`Source` node points to a single OpenCL source file.
|
||||
|
||||
| Attribute Name | \# |Description|
|
||||
|-----|-----|-----|
|
||||
| `filename` | (1) | Name of the file containing OpenCL source code. Note that the path is relative to your executable. Multiple source nodes will have their sources concatenated in order. |
|
||||
|
||||
**Sub-nodes**: None
|
||||
|
||||
### Define Node and Sub-Node Structure
|
||||
|
||||
`Define` node configures a single `#‍define` instruction to be added to
|
||||
the sources during compilation (JIT).
|
||||
|
||||
| Attribute Name | \# | Description |
|
||||
|------|-------|------|
|
||||
| `name` | (1) | The name of the defined JIT. For static constants, this can include the value as well, which is taken as a string. |
|
||||
| `param` | (0/1) | This parameter value is used as the value of this JIT definition. |
|
||||
| `type` | (0/1) | The parameter type. Accepted values: `int`, `float`, and `int[]`, `float[]` for arrays. |
|
||||
| `default` | (0/1) | The default value to be used if the specified parameters are missing from the operation in the IR. |
|
||||
|
||||
**Sub-nodes:** None
|
||||
|
||||
The resulting JIT has the following form:
|
||||
`#‍define [name] [type] [value/default]`.
|
||||
|
||||
### Buffers Node and Sub-Node Structure
|
||||
|
||||
`Buffers` node configures all input/output buffers for the OpenCL entry
|
||||
function. No buffers node structure exists.
|
||||
|
||||
**Sub-nodes:** `Data` (0+), `Tensor` (1+)
|
||||
|
||||
### Data Node and Sub-Node Structure
|
||||
|
||||
`Data` node configures a single input with static data, for example,
|
||||
weights or biases.
|
||||
|
||||
| Attribute Name | \# | Description |
|
||||
|----|-----|------|
|
||||
| `name` | (1) | Name of a blob attached to an operation in the IR |
|
||||
| `arg-index` | (1) | 0-based index in the entry function arguments to be bound to |
|
||||
|
||||
**Sub-nodes**: None
|
||||
|
||||
### Tensor Node and Sub-Node Structure
|
||||
|
||||
`Tensor` node configures a single input or output tensor.
|
||||
|
||||
| Attribute Name | \# | Description |
|
||||
|------|-------|-------|
|
||||
| `arg-index` | (1) | 0-based index in the entry function arguments to be bound to. |
|
||||
| `type` | (1) | `input` or `output` |
|
||||
| `port-index` | (1) | 0-based index in the operation input/output ports in the IR |
|
||||
| `format` | (0/1) | Data layout declaration for the tensor. Accepted values: `BFYX`, `BYXF`, `YXFB`, `FYXB`, and same values in all lowercase. Default value: `BFYX` |
|
||||
|
||||
### CompilerOptions Node and Sub-Node Structure
|
||||
|
||||
`CompilerOptions` node configures the compilation flags for the OpenCL
|
||||
sources.
|
||||
|
||||
| Attribute Name | \# | Description |
|
||||
|--------|-----|------|
|
||||
| `options` | (1) | Options string to be passed to the OpenCL compiler |
|
||||
|
||||
**Sub-nodes**: None
|
||||
|
||||
### WorkSizes Node and Sub-Node Structure
|
||||
|
||||
`WorkSizes` node configures the global/local work sizes to be used when
|
||||
queuing an OpenCL program for execution.
|
||||
|
||||
| Attribute Name | \# | Description |
|
||||
|-----|------|-----|
|
||||
| `global`<br>`local` | (0/1)<br>(0/1) | An array of up to three integers or formulas for defining OpenCL work-sizes to be used during execution.<br> The formulas can use the values of the B,F,Y,X dimensions and contain the operators: +,-,/,\*,%. All operators are evaluated in integer arithmetic. <br>Default value: `global=”B*F*Y*X” local=””` |
|
||||
| `dim` | (0/1) | A tensor to take the work-size from. Accepted values: `input N`, `output`, where `N` is an index of input tensor starting with 0. Default value: `output` |
|
||||
|
||||
**Sub-nodes**: None
|
||||
|
||||
## Example Configuration File
|
||||
|
||||
The following code sample provides an example configuration file in XML
|
||||
format. For information on the configuration file structure, see
|
||||
[Configuration File Format](#config-file-format).
|
||||
```xml
|
||||
<CustomLayer name="ReLU" type="SimpleGPU" version="1">
|
||||
<Kernel entry="example_relu_kernel">
|
||||
<Source filename="custom_layer_kernel.cl"/>
|
||||
<Define name="neg_slope" type="float" param="negative_slope" default="0.0"/>
|
||||
</Kernel>
|
||||
<Buffers>
|
||||
<Tensor arg-index="0" type="input" port-index="0" format="BFYX"/>
|
||||
<Tensor arg-index="1" type="output" port-index="0" format="BFYX"/>
|
||||
</Buffers>
|
||||
<CompilerOptions options="-cl-mad-enable"/>
|
||||
<WorkSizes global="X,Y,B*F"/>
|
||||
</CustomLayer>
|
||||
```
|
||||
|
||||
## Built-In Definitions for Custom Layers
|
||||
|
||||
The following table includes definitions that are attached before
|
||||
user sources, where `<TENSOR>` is the actual input and output, for
|
||||
example, `INPUT0` or `OUTPUT0`.
|
||||
|
||||
For an example, see [Example Kernel](#example-kernel).
|
||||
|
||||
| Name | Value |
|
||||
|---|---|
|
||||
| `NUM_INPUTS` | Number of the input tensors bound to this kernel |
|
||||
| `GLOBAL_WORKSIZE` | An array of global work sizes used to execute this kernel |
|
||||
| `GLOBAL_WORKSIZE_SIZE` | The size of the `GLOBAL_WORKSIZE` array |
|
||||
| `LOCAL_WORKSIZE` | An array of local work sizes used to execute this kernel |
|
||||
| `LOCAL_WORKSIZE_SIZE` | The size of the `LOCAL_WORKSIZE` array |
|
||||
| `<TENSOR>_DIMS`| An array of the tensor dimension sizes. Always ordered as `BFYX` |
|
||||
| `<TENSOR>_DIMS_SIZE`| The size of the `<TENSOR>_DIMS` array.|
|
||||
| `<TENSOR>_TYPE`| The datatype of the tensor: `float`, `half`, or `char`|
|
||||
| `<TENSOR>_FORMAT_` | The format of the tensor, BFYX, BYXF, YXFB , FYXB, or ANY. The format is concatenated to the defined name. You can use the tensor format to define codepaths in your code with `#‍ifdef/#‍endif`. |
|
||||
| `<TENSOR>_LOWER_PADDING` | An array of padding elements used for the tensor dimensions before they start. Always ordered as BFYX.|
|
||||
| `<TENSOR>_ LOWER_PADDING_SIZE` | The size of the `<TENSOR>_LOWER_PADDING` array |
|
||||
| `<TENSOR>_UPPER_PADDING` | An array of padding elements used for the tensor dimensions after they end. Always ordered as BFYX. |
|
||||
| `<TENSOR>_UPPER_PADDING_SIZE` | The size of the `<TENSOR>_UPPER_PADDING` array |
|
||||
| `<TENSOR>_PITCHES` | The number of elements between adjacent elements in each dimension. Always ordered as BFYX.|
|
||||
| `<TENSOR>_PITCHES_SIZE`| The size of the `<TENSOR>_PITCHES` array |
|
||||
| `<TENSOR>_OFFSET`| The number of elements from the start of the tensor to the first valid element, bypassing the lower padding. |
|
||||
All `<TENSOR>` values are automatically defined for every tensor
|
||||
bound to this operation, such as `INPUT0`, `INPUT1`, and `OUTPUT0`, as shown
|
||||
in the following example:
|
||||
|
||||
```sh
|
||||
#define INPUT0_DIMS_SIZE 4
|
||||
#define INPUT0_DIMS (int []){ 1,96,55,55, }
|
||||
```
|
||||
|
||||
## Example Kernel<a name="example-kernel"></a>
|
||||
|
||||
```c
|
||||
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
|
||||
__kernel void example_relu_kernel(
|
||||
const __global INPUT0_TYPE* input0,
|
||||
__global OUTPUT0_TYPE* output)
|
||||
{
|
||||
const uint idx = get_global_id(0);
|
||||
const uint idy = get_global_id(1);
|
||||
const uint idbf = get_global_id(2);//batches*features, as OpenCL supports 3D nd-ranges only
|
||||
const uint feature = idbf%OUTPUT0_DIMS[1];
|
||||
const uint batch = idbf/OUTPUT0_DIMS[1];
|
||||
//notice that pitches are in elements, not in bytes!
|
||||
const uint in_id = batch*INPUT0_PITCHES[0] + feature*INPUT0_PITCHES[1] + idy*INPUT0_PITCHES[2] + idx*INPUT0_PITCHES[3] + INPUT0_OFFSET;
|
||||
const uint out_id = batch*OUTPUT0_PITCHES[0] + feature*OUTPUT0_PITCHES[1] + idy*OUTPUT0_PITCHES[2] + idx*OUTPUT0_PITCHES[3] + OUTPUT0_OFFSET;
|
||||
|
||||
INPUT0_TYPE value = input0[in_id];
|
||||
//neg_slope (which is non-zero for leaky ReLU) is put automatically as #define, refer to the config xml
|
||||
output[out_id] = value < 0 ? value * neg_slope : value;
|
||||
}
|
||||
```
|
||||
|
||||
> **NOTE**: As described in the previous section, all items like
|
||||
> `INPUT0_TYPE` are actually defined as OpenCL (pre-)compiler inputs by
|
||||
> the Inference Engine for efficiency reasons. See [Debugging
|
||||
> Tips](#debugging-tips) for information on debugging the results.
|
||||
|
||||
> **NOTE**: Several GPU-targeted kernels are also added to the binaries upon compilation of samples
|
||||
> so that the sample application can easy load them.
|
||||
> Refer to the `cldnn_global_custom_kernels` folder in the GPU plugin installation directory.
|
||||
|
||||
## Debugging Tips<a name="debugging-tips"></a>
|
||||
|
||||
* **Using `printf` in the OpenCL™ Kernels**.
|
||||
To debug the specific values, you can use `printf` in your kernels.
|
||||
However, be careful not to output excessively, which
|
||||
could generate too much data. The `printf` output is typical, so
|
||||
your output can be truncated to fit the buffer. Also, because of
|
||||
buffering, you actually get an entire buffer of output when the
|
||||
execution ends.<br>
|
||||
|
||||
For more information, refer to the [printf
|
||||
Function](https://www.khronos.org/registry/OpenCL/sdk/1.2/docs/man/xhtml/printfFunction.html).
|
||||
@@ -1,60 +0,0 @@
|
||||
# Inference Engine Extensibility Mechanism {#openvino_docs_IE_DG_Extensibility_DG_Intro}
|
||||
|
||||
@sphinxdirective
|
||||
|
||||
.. toctree::
|
||||
:maxdepth: 1
|
||||
:hidden:
|
||||
|
||||
openvino_docs_IE_DG_Extensibility_DG_AddingNGraphOps
|
||||
openvino_docs_IE_DG_Extensibility_DG_Custom_ONNX_Ops
|
||||
CPU Kernels Extensibility <openvino_docs_IE_DG_Extensibility_DG_CPU_Kernel>
|
||||
GPU Kernels Extensibility <openvino_docs_IE_DG_Extensibility_DG_GPU_Kernel>
|
||||
VPU Kernels Extensibility <openvino_docs_IE_DG_Extensibility_DG_VPU_Kernel>
|
||||
openvino_docs_IE_DG_Extensibility_DG_Extension
|
||||
openvino_docs_IE_DG_Extensibility_DG_Building
|
||||
|
||||
@endsphinxdirective
|
||||
|
||||
If your model contains operations not normally supported by OpenVINO, the Inference Engine Extensibility API lets you add support for those custom operations in a library containing custom nGraph operation sets, corresponding extensions to the Model Optimizer, and a device plugin extension. See the overview in the [Custom Operations Guide](../../HOWTO/Custom_Layers_Guide.md) to learn how these work together.
|
||||
|
||||
To load the Extensibility library to the `InferenceEngine::Core` object, use the `InferenceEngine::Core::AddExtension` method.
|
||||
|
||||
## Inference Engine Extension Library
|
||||
|
||||
An Inference Engine Extension dynamic library contains the following components:
|
||||
|
||||
* [Extension Library](Extension.md):
|
||||
- Contains custom operation sets
|
||||
- Provides CPU implementations for custom operations
|
||||
* [Custom nGraph Operation](AddingNGraphOps.md):
|
||||
- Enables the use of `InferenceEngine::Core::ReadNetwork` to read Intermediate Representation (IR) with unsupported
|
||||
operations
|
||||
- Enables the creation of `ngraph::Function` with unsupported operations
|
||||
- Provides a shape inference mechanism for custom operations
|
||||
|
||||
> **NOTE**: This documentation is written based on the [Template extension](https://github.com/openvinotoolkit/openvino/tree/master/docs/template_extension), which demonstrates extension development details. You can review the complete code, which is fully compilable and up-to-date, to see how it works.
|
||||
|
||||
## Execution Kernels
|
||||
|
||||
The Inference Engine workflow involves the creation of custom kernels and either custom or existing operations.
|
||||
|
||||
An _operation_ is a network building block implemented in the training framework, for example, `Convolution` in Caffe*.
|
||||
A _kernel_ is defined as the corresponding implementation in the Inference Engine.
|
||||
|
||||
Refer to the [Model Optimizer Extensibility](../../MO_DG/prepare_model/customize_model_optimizer/Customize_Model_Optimizer.md)
|
||||
for details on how a mapping between framework operations and Inference Engine kernels is registered.
|
||||
|
||||
In short, you can plug your own kernel implementations into the Inference Engine and map them to the operations in the original framework.
|
||||
|
||||
The following pages describe how to integrate custom _kernels_ into the Inference Engine:
|
||||
|
||||
* [Introduction to development of custom CPU kernels](CPU_Kernel.md)
|
||||
* [Introduction to development of custom GPU kernels](GPU_Kernel.md)
|
||||
* [Introduction to development of custom VPU kernels](VPU_Kernel.md)
|
||||
|
||||
## See Also
|
||||
|
||||
* [Build an extension library using CMake*](Building.md)
|
||||
* [Using Inference Engine Samples](../Samples_Overview.md)
|
||||
* [Hello Shape Infer SSD sample](../../../samples/cpp/hello_reshape_ssd/README.md)
|
||||
@@ -1,682 +0,0 @@
|
||||
# How to Implement Custom Layers for VPU (Intel® Neural Compute Stick 2) {#openvino_docs_IE_DG_Extensibility_DG_VPU_Kernel}
|
||||
|
||||
To enable operations not supported by OpenVINO™ out of the box, you need a custom extension for Model Optimizer, a custom nGraph operation set, and a custom kernel for the device you will target. This page describes custom kernel support for one the VPU, the Intel® Neural Compute Stick 2 device, which uses the MYRIAD device plugin.
|
||||
|
||||
> **NOTES:**
|
||||
> * OpenCL\* custom layer support is available in the preview mode.
|
||||
> * This section assumes you are familiar with developing kernels using OpenCL.
|
||||
|
||||
To customize your topology with an OpenCL layer, carry out the tasks described on this page:
|
||||
|
||||
1. Write and compile your OpenCL code with the standalone offline OpenCL compiler (`clc`).
|
||||
2. Write a configuration file to bind the OpenCL kernel to the topology file (`.xml`) of the model IR.
|
||||
3. Pass the configuration file to the Inference Engine with the model IR.
|
||||
|
||||
## Compile OpenCL code for VPU (Intel® Neural Compute Stick 2)
|
||||
|
||||
> **NOTE**: OpenCL compiler, targeting Intel® Neural Compute Stick 2 for the SHAVE* processor only, is redistributed with OpenVINO.
|
||||
OpenCL support is provided by ComputeAorta* and is distributed under a license agreement between Intel® and Codeplay* Software Ltd.
|
||||
|
||||
The OpenCL toolchain for the Intel® Neural Compute Stick 2 supports offline compilation only, so first compile OpenCL C code using the standalone `clc` compiler. You can find the compiler binary at `<INSTALL_DIR>/tools/cl_compiler`.
|
||||
|
||||
> **NOTE**: By design, custom OpenCL layers support any OpenCL kernels written assuming OpenCL version 1.2. It also supports half float extension and is optimized for this type, because it is a native type for Intel® Movidius™ VPUs.
|
||||
|
||||
1. Prior to running a compilation, make sure that the following variables are set:
|
||||
* `SHAVE_MA2X8XLIBS_DIR=<INSTALL_DIR>/tools/cl_compiler/lib/`
|
||||
* `SHAVE_LDSCRIPT_DIR=<INSTALL_DIR>/tools/cl_compiler/ldscripts/`
|
||||
* `SHAVE_MYRIAD_LD_DIR=<INSTALL_DIR>/tools/cl_compiler/bin/`
|
||||
* `SHAVE_MOVIASM_DIR=<INSTALL_DIR>/tools/cl_compiler/bin/`
|
||||
2. Run the compilation with the command below. You should use `--strip-binary-header` to make an OpenCL runtime-agnostic binary runnable with the Inference Engine.
|
||||
```bash
|
||||
cd <INSTALL_DIR>/tools/cl_compiler/bin
|
||||
./clc --strip-binary-header custom_layer.cl -o custom_layer.bin
|
||||
```
|
||||
|
||||
## Write a Configuration File
|
||||
|
||||
To tie the topology IR for a layer you customize, prepare a configuration file, so that the Inference Engine can find parameters for your kernel and the execution work grid is described.
|
||||
For example, consider the following OpenCL kernel signature:
|
||||
```cpp
|
||||
__kernel void reorg_nhwc(__global const half *src, __global half *out, int w, int h, int c, int stride);
|
||||
```
|
||||
A configuration file for this kernel might be the following:
|
||||
```xml
|
||||
<CustomLayer name="ReorgYolo" type="MVCL" version="1">
|
||||
<Kernel entry="reorg_nhwc">
|
||||
<Source filename="reorg.bin"/>
|
||||
</Kernel>
|
||||
<Parameters>
|
||||
<Tensor arg-name="src" type="input" port-index="0" format="BYXF"/>
|
||||
<Tensor arg-name="out" type="output" port-index="0" format="BYXF"/>
|
||||
<Scalar arg-name="w" type="int" port-index="0" source="I.X" />
|
||||
<Scalar arg-name="h" type="int" port-index="0" source="I.Y" />
|
||||
<Scalar arg-name="c" type="int" port-index="0" source="I.F" />
|
||||
<Scalar arg-name="stride" type="int" source="stride" />
|
||||
</Parameters>
|
||||
<WorkSizes dim="input,0" global="(Y+7)/8*8,1,1" local="8,1,1"/>
|
||||
</CustomLayer>
|
||||
```
|
||||
Each custom layer is described with the `CustomLayer` node. It has the following nodes and attributes:
|
||||
- Root node `CustomLayer` contains the following attributes:
|
||||
- `name` – (Required) The name of the Inference Engine layer to bind the kernel with.
|
||||
- `type` and `version` – (Required) Reserved for future use. Set them to `MVCL` and `1` respectively.
|
||||
- `max-shaves` – (Optional) The maximum number of SHAVE cores that should be dedicated for the layer. It is useful for debugging concurrency issues or for resource saving that memory bound kernel does not scale well with the number of cores, so more resources can be left for the rest of a topology.
|
||||
- Sub-node `Kernel` must contain the following attributes:
|
||||
- `entry` – The name of your kernel function as you defined it in a source file. In the example above, it is `reorg_nhwc`.
|
||||
- Node `Source` must contain the following attributes:
|
||||
- `filename` – The path to a compiled binary relative to the XML configuration file.
|
||||
- Sub-node `Parameters` – Describes parameters bindings. For more information, see the description below.
|
||||
- Sub-node `WorkSizes` – Describes local and global work group sizes and the source for dimension deduction as a pair `direction,port`. In the example above, the work group is described relatively to the dimension of the input tensor that comes through port 0 in the IR. `global` and `local` work group configurations support any simple math expressions with +,-,\*,/, and () from `B`(batch), `Y`(height), `X`(width) and `F`(channels).
|
||||
- Sub-node `Where` – Allows to customize bindings with the `key="value"` attribute. For example, to substitute only 3x3 convolutions, write `<Where kernel="3,3"/>` in the binding xml.
|
||||
|
||||
Parameter description supports `Tensor` of one of tensor types such as `input`, `output`, `input_buffer`, `output_buffer` or `data`, `Scalar`, or `Data` nodes and has the following format:
|
||||
- Each `Tensor` node of `input` or `output` type must contain the following attributes:
|
||||
- `arg-name` – The name of a kernel parameter in the kernel signature.
|
||||
- `type` – Node type: `input` or `output` as specified in the IR.
|
||||
- `port-index` – A number of input/output ports as specified in the IR.
|
||||
- `format` – The channel order in the tensor. Optional conversion layers are generated if the custom layer format is not compatible with formats of neighboring layers. `BFXY`, `BYXF`, and `ANY` formats are supported currently.
|
||||
- Each `Tensor` node of `input_buffer` or `output_buffer` type must contain the following attributes:
|
||||
- `arg-name` – The name of a kernel parameter in the kernel signature.
|
||||
- `type` – Node type: `input_buffer` or `output_buffer`. Use the appropriate type to bind multiple kernels that correspond to different stages of the same layer.
|
||||
- `port-index` – The unique identifier to bind by.
|
||||
- `dim` – The dim source with the same `direction,port` format used for `WorkSizes` bindings.
|
||||
- `size` – Amount of bytes needed. Current expression syntax supports only expression over dimensions of over selected input/output tensor or constants and might be expended in the future.
|
||||
|
||||
Here is an example of multi-stage MVN layer binding:
|
||||
```xml
|
||||
<CustomLayer name="MVN" stage="0" type="MVCL" version="1">
|
||||
<Kernel entry="reduction_mean">
|
||||
<Source filename="mvn.bin"/>
|
||||
</Kernel>
|
||||
<Parameters>
|
||||
<Tensor arg-name="src" type="input" port-index="0" format="BFYX"/>
|
||||
<Tensor arg-name="mean" type="output_buffer" port-index="0" dim="output,0" size="Y*F*4"/>
|
||||
<Tensor arg-name="variance" type="output_buffer" port-index="1" dim="output,0" size="Y*F*4"/>
|
||||
<!--other parameters -->
|
||||
</Parameters>
|
||||
<WorkSizes dim="output,0" global="((Y+7)/8)*8,F,1" local="8,1,1"/>
|
||||
</CustomLayer>
|
||||
<CustomLayer name="MVN" stage="1" type="MVCL" version="1">
|
||||
<Kernel entry="mvn_scale">
|
||||
<Source filename="mvn_scale_changed_orded.bin"/>
|
||||
</Kernel>
|
||||
<Parameters>
|
||||
<Tensor arg-name="src_data" type="input" port-index="0" format="BFYX"/>
|
||||
<Tensor arg-name="dst_data" type="output" port-index="0" format="BFYX"/>
|
||||
<Tensor arg-name="mean_part" type="input_buffer" port-index="0" dim="output,0" size="Y*F*4"/>
|
||||
<Tensor arg-name="power_mean" type="input_buffer" port-index="1" dim="output,0" size="Y*F*4"/>
|
||||
<!--other parameters -->
|
||||
</Parameters>
|
||||
<WorkSizes dim="output,0" global="((Y+7)/8)*8,F,1" local="8,1,1"/>
|
||||
</CustomLayer>
|
||||
```
|
||||
- Each `Tensor` node that has the type `data` must contain the following attributes:
|
||||
- `source` – A name of the blob as it is in the IR. Typical example is `weights` for convolution.
|
||||
- `format` – Specifies the channel order in the tensor. Optional conversion layers are generated if the custom layer format is not.
|
||||
```xml
|
||||
<CustomLayer name="BinaryConvolution" type="MVCL" version="1">
|
||||
<Kernel entry="binary_convolution">
|
||||
<Source filename="binary_layers.bin"/>
|
||||
</Kernel>
|
||||
<Parameters>
|
||||
<Tensor arg-name="src_data" type="input" port-index="0" format="BFYX"/>
|
||||
<Data arg-name="weights_data" type="data" source="weights" format="ANY"/>
|
||||
<Tensor arg-name="dst_data" type="output" port-index="0" format="BFYX"/>
|
||||
<!--other parameters -->
|
||||
</Parameters>
|
||||
<WorkSizes dim="output,0" global="X,Y,F" local="1,1,1"/>
|
||||
</CustomLayer>
|
||||
```
|
||||
- Each `Scalar` node must contain the following attributes:
|
||||
- `arg-name` – The name of a kernel parameter in the kernel signature.
|
||||
- `type` – `int` or `float` value. It is used for correct argument extraction from IR parameters.
|
||||
- `source` – Contains the name of the parameter in the IR file or input/output (`I`/`O`, `In`/`On`, where `n` is a port number)
|
||||
followed by dimension `B`(batch), `Y`(height), `X`(width), or `F`(channels).
|
||||
|
||||
- Each `Data` node must contain the following attributes:
|
||||
- `arg-name` – The name of a kernel parameter in the kernel signature.
|
||||
- `type` – Node type. Currently, `local_data` is the only supported value, which defines buffer allocated in fast local on-chip memory. It is limited to 100KB for all `__local` and
|
||||
`__private` arrays defined inside the kernel as well as all `__local` parameters passed to the kernel. Note that a manual-DMA extension requires double buffering.
|
||||
If the custom layer is detected to run out of local memory, the inference fails.
|
||||
- `dim` – The dim source with the same `direction,port` format used for `WorkSizes` bindings.
|
||||
- `size` – Amount of bytes needed. The current expression syntax supports only expression over dimensions of over selected input/output tensor or constants and may be extended in the future.
|
||||
The example binding below illustrates a kernel with two local buffers passed to the kernel.
|
||||
```xml
|
||||
<CustomLayer name="GRN" type="MVCL" version="1">
|
||||
<Kernel entry="grn_NCHW">
|
||||
<Source filename="grn.bin"/>
|
||||
</Kernel>
|
||||
<Parameters>
|
||||
<Tensor arg-name="src_data" type="input" port-index="0" format="BFYX"/>
|
||||
<Tensor arg-name="dst_data" type="output" port-index="0" format="BFYX"/>
|
||||
<Data arg-name="src" type="local_data" dim="input,0" size="X*F*2" />
|
||||
<Data arg-name="dst" type="local_data" dim="input,0" size="X*F*2" />
|
||||
<Scalar arg-name="C" type="int" port-index="0" source="I.F" />
|
||||
<Scalar arg-name="bias" type="float" source="bias" />
|
||||
</Parameters>
|
||||
<WorkSizes dim="input,0" global="X,Y,1" local="X,1,1"/>
|
||||
</CustomLayer>
|
||||
```
|
||||
|
||||
## Pass Configuration File to Inference Runtime
|
||||
|
||||
> **NOTE**: If both native and custom layer implementations are present, the custom kernel has a priority over the native one.
|
||||
|
||||
Before loading the network that features the custom layers, provide a separate configuration file and load it using the InferenceEngine::Core::SetConfig() method with the PluginConfigParams::KEY_CONFIG_FILE key and the configuration file name as a value:
|
||||
```cpp
|
||||
InferenceEngine::Core core;
|
||||
// Load custom layers
|
||||
core.SetConfig({ { InferenceEngine::PluginConfigParams::KEY_CONFIG_FILE, "<path to the xml file>" } }, "MYRIAD");
|
||||
```
|
||||
Optionally, set a path to a custom layers description with a pair of `VPU_CUSTOM_LAYERS` and `/path/to/your/customLayers.xml`
|
||||
as a network configuration:
|
||||
```cpp
|
||||
InferenceEngine::Core core;
|
||||
std::map<std::string, std::string> networkConfig;
|
||||
config["VPU_CUSTOM_LAYERS"] = "/path/to/your/customLayers.xml";
|
||||
// Load custom layers in network config
|
||||
auto exeNetwork = core.LoadNetwork(cnnNetwork, "MYRIAD", networkConfig);
|
||||
```
|
||||
|
||||
## Optimizing Kernels with OpenCL for VPU (Intel® Neural Compute Stick 2)
|
||||
|
||||
This section provides optimization guidelines on writing custom layers with OpenCL for VPU devices. Knowledge about general OpenCL
|
||||
programming model and OpenCL kernel language is assumed and not a subject of this section. The OpenCL model mapping to VPU is described in the table below.
|
||||
|
||||
| OpenCL Model | VPU Mapping|
|
||||
|-----|----|
|
||||
| Device code | Executed on SHAVE cores |
|
||||
| Private memory | Mapped to CMX internal memory, limited to 100KB per work group, valid only while the work group is executed |
|
||||
| Local memory | Mapped to CMX internal memory, limited to 100KB per work group, valid only while the work group is executed |
|
||||
| Global memory | Mapped to DDR, used to pass execution preserved parameters for inputs, outputs, and blobs |
|
||||
| Work group | Executed on a single SHAVE core iterating over multiple work items |
|
||||
|
||||
Note that by the OpenCL specification, the work group execution order is not specified. This means that it is your
|
||||
responsibility to ensure that race conditions among work groups are not introduced. Custom layer runtime spits evenly
|
||||
work grid among available compute resources and executes them in an arbitrary order. This static scheduling approach works best if the load is evenly spread out across work groups, which is a typical case for Deep Learning kernels. The following guidelines are recommended to use for work group partitioning:
|
||||
|
||||
1. Split work evenly across work groups.
|
||||
2. Adjust work group granularity to maintain equal workload for all compute codes.
|
||||
3. Set the maximum number of cores using the `max-shaves` attribute for the `CustomLayer` node. This keeps more resources for the rest of topology. It is also useful if the kernel scalability reached its limits, which may happen while optimizing memory bound kernels or kernels with poor parallelization.
|
||||
4. Try an alternate data layout (`BFXY`/`BYXF`) for the kernel if it improves work group partitioning or data access patterns.
|
||||
Consider not just specific layer boost, but full topology performance because data conversion layers would be automatically inserted
|
||||
as appropriate.
|
||||
|
||||
Offline OpenCL compiler (`clc`) features automatic vectorization over `get_global_id(0)` usage, if uniform access is detected.
|
||||
For example, the kernel below could be automatically vectorized:
|
||||
```cpp
|
||||
__kernel void cvtf32f16(__global float* restrict inImage, __global half* restrict outImage,
|
||||
float scale, float bais)
|
||||
{
|
||||
int idx = get_global_id(0) + get_global_id(1) * get_global_size(0) + get_global_id(2) * get_global_size(0) * get_global_size(1);
|
||||
outImage[idx] = convert_half(inImage[idx]*scale+bais);
|
||||
}
|
||||
```
|
||||
However, this work-group based vectorizer (WGV) conflicts with the default LLVM vectorizer based on superword level parallelism
|
||||
(SLP) for the current compiler version. Manual vectorization is recommended to provide the best performance for non-uniform code
|
||||
patterns. WGV works if and only if vector types are not used in the code.
|
||||
|
||||
Here is a short list of optimization tips:
|
||||
|
||||
1. Help auto-vectorizer ensure non-aliasing pointers for kernel parameters by putting `restrict` where possible.
|
||||
- This can give a performance boost, especially for kernels with unrolling, like `ocl_grn` from the example below.
|
||||
- Place `restrict` markers for kernels with manually vectorized codes. In the `ocl_grn` kernel below, the unrolled version without `restrict` is up to 20% slower than the most optimal one, which combines unrolling and `restrict`.
|
||||
2. Put `#‍pragma unroll N` to your loop header. The compiler does not trigger unrolling by default, so it is your responsibility to
|
||||
annotate the code with pragmas as appropriate. The `ocl_grn` version with `#‍pragma unroll 4` is up to 50% faster, most of which comes from unrolling the first loop, because LLVM, in general, is better in scheduling 3-stage loops (load-compute-store), while the fist loop
|
||||
`variance += (float)(src_data[c*H*W + y*W + x] * src_data[c*H*W + y*W + x]);` is only 2-stage (load-compute). Pay
|
||||
attention to unrolling such cases first. Unrolling factor is loop-dependent. Choose the smallest number that
|
||||
still improves performance as an optimum between the kernel size and execution speed. For this specific kernel, changing the unroll factor from `4` to `6` results in the same performance, so unrolling factor equal to 4 is an optimum. For Intel® Neural Compute Stick 2, unrolling is conjugated with the automatic software pipelining for load, store, and compute stages:
|
||||
```cpp
|
||||
__kernel void ocl_grn(__global const half* restrict src_data, __global half* restrict dst_data, int C, float bias)
|
||||
{
|
||||
int x = get_global_id(0);
|
||||
int W = get_global_size(0);
|
||||
int y = get_global_id(1);
|
||||
int H = get_global_size(1);
|
||||
|
||||
float variance = bias + 1e-9f;
|
||||
|
||||
#pragma unroll 4
|
||||
for (int c = 0; c < C; c++)
|
||||
variance += (float)(src_data[c*H*W + y*W + x] * src_data[c*H*W + y*W + x]);
|
||||
|
||||
variance = 1.f / native_sqrt(variance);
|
||||
|
||||
#pragma unroll 4
|
||||
for (int c = 0; c < C; c++)
|
||||
dst_data[c*H*W + y*W + x] = (half)((float)src_data[c*H*W + y*W + x] * variance);
|
||||
}
|
||||
```
|
||||
To check the efficiency of WGV, you can compare performance of the kernel above with the kernel below, which is manually vectorized over width:
|
||||
```cpp
|
||||
__kernel void ocl_grn_line(__global const half* restrict src_data, __global half* restrict dst_data, int C, int W, float bias)
|
||||
{
|
||||
int y = get_global_id(1);
|
||||
int H = get_global_size(1);
|
||||
|
||||
for (int x = 0; x < W/8; x++)
|
||||
{
|
||||
float8 variance = (float8)(bias+1e-9f);
|
||||
|
||||
#pragma unroll 4
|
||||
for (int c = 0; c < C; c++)
|
||||
{
|
||||
__global const half8* restrict src_line = ((__global const half8 * restrict)(src_data + c*H*W + y*W));
|
||||
half8 sh = src_line[x];
|
||||
variance += convert_float8(sh*sh);
|
||||
}
|
||||
|
||||
variance = 1.f/native_sqrt(variance);
|
||||
|
||||
#pragma unroll 4
|
||||
for (int c = 0; c < C; c++)
|
||||
{
|
||||
__global const half8* restrict src_line = ((__global const half8 * restrict)(src_data + c*H*W + y*W));
|
||||
__global half8* restrict dst_line = ((__global half8 * restrict)(dst_data + c*H*W + y*W));
|
||||
|
||||
dst_line[x] = convert_half8(convert_float8(src_line[x])*variance);
|
||||
}
|
||||
}
|
||||
for (int x = W/8*8; x < W; x++)
|
||||
{
|
||||
float variance = bias+1e-9f;
|
||||
#pragma unroll 4
|
||||
for (int c = 0; c < C; c++)
|
||||
variance += (float)(src_data[c*H*W + y*W + x]*src_data[c*H*W + y*W + x]);
|
||||
|
||||
variance = 1.f/native_sqrt(variance);
|
||||
|
||||
#pragma unroll 4
|
||||
for (int c = 0; c < C; c++)
|
||||
dst_data[c*H*W + y*W + x] = (float)src_data[c*H*W + y*W + x]*variance;
|
||||
}
|
||||
}
|
||||
```
|
||||
Both versions perform the same, but the second one has more complex code.
|
||||
|
||||
3. If it is easy to predict the work group size, you can also use the `reqd_work_group_size` kernel attribute to ask the compiler
|
||||
to unroll the code up to the local size of the work group. Note that if the kernel is actually executed with the
|
||||
different work group configuration, the result is undefined.
|
||||
|
||||
4. Prefer to use the `half` compute if it keeps reasonable accuracy. 16-bit float is a native type for Intel® Neural Compute Stick 2, most of the functions `half_*` are mapped to a single hardware instruction.
|
||||
Use the standard `native_*` function for the rest of types.
|
||||
|
||||
5. Prefer to use the `convert_half` function over `vstore_half` if conversion to 32-bit float is required. `convert_half` is mapped to a single hardware instruction. For the `cvtf32f16` kernel above, the line `outImage[idx] = convert_half(inImage[idx]*scale+bais);` is eight times slower than the code with `vstore_half`.
|
||||
|
||||
6. Mind early exits. Early exit can be extremely costly for the current version of the `clc` compiler due to conflicts with the
|
||||
auto-vectorizer. The generic advice would be to setup local size by `x` dimension equal to inputs or/and outputs width.
|
||||
If it is impossible to define the work grid that exactly matches inputs or/and outputs to eliminate checks, for example,
|
||||
`if (get_global_id(0) >= width) return`, use line-wise kernel variant with manual vectorization.
|
||||
The kernel example below demonstrates the impact of early exits on kernel performance.
|
||||
```cpp
|
||||
// Initial version
|
||||
__kernel void reorg(const __global half* restrict src, __global half* restrict out, int stride)
|
||||
{
|
||||
int w = get_global_id(0);
|
||||
int W = get_global_size(0);
|
||||
|
||||
int h = get_global_id(1);
|
||||
int H = get_global_size(1);
|
||||
|
||||
int c = get_global_id(2);
|
||||
int C = get_global_size(2);
|
||||
|
||||
int C2 = C/(stride*stride);
|
||||
int offset = c / C2;
|
||||
int c2 = c - C2 * offset;
|
||||
|
||||
int H2 = H*stride;
|
||||
int W2 = W*stride;
|
||||
|
||||
int h2 = h*stride + offset / stride;
|
||||
int w2 = w*stride + offset - stride * (offset / stride);
|
||||
|
||||
out[W*H*c + W*h + w] = src[W2*H2*c2 + W2*h2 + w2];
|
||||
}
|
||||
```
|
||||
This `reorg` kernel is auto-vectorizable, but an input for YOLO v2 topology is `NCHW=<1,64,26,26>` and it is not multiple of vector width, which is `8` for `half` data type. As a result, the Inference Engine does not select the auto-vectorized kernel.
|
||||
To compare performance of auto-vectorized and scalar version of the kernel, change the input size to`NCHW=<1,64,26,32>`. This enables the auto-vectorized version to be selected by the Inference Engine and can give you about 30% uplift.
|
||||
Since the auto-vectorized version is faster, it makes sense to enable it for the YOLO v2 topology input size by setting the local size multiple of vector, for example, 32, and adjust global sizes accordingly. As a result, the execution work grid exceeds actual input dimension, so out-of-bound checks should be inserted. See the updated kernel version below:
|
||||
```cpp
|
||||
// Version with out-of-bound checks added
|
||||
__kernel void reorg(const __global half* restrict src, __global half* restrict out, int W, int stride)
|
||||
{
|
||||
int w = get_global_id(0);
|
||||
w = min(w, W-1);
|
||||
|
||||
int h = get_global_id(1);
|
||||
int H = get_global_size(1);
|
||||
|
||||
int c = get_global_id(2);
|
||||
int C = get_global_size(2);
|
||||
|
||||
int C2 = C/(stride*stride);
|
||||
int offset = c / C2;
|
||||
int c2 = c - C2 * offset;
|
||||
|
||||
int H2 = H*stride;
|
||||
int W2 = W*stride;
|
||||
|
||||
int h2 = h*stride + offset / stride;
|
||||
int w2 = w*stride + offset - stride * (offset / stride);
|
||||
|
||||
out[W*H*c + W*h + w] = src[W2*H2*c2 + W2*h2 + w2];
|
||||
}
|
||||
```
|
||||
This code performs the same as the initial kernel above (scalar) due to branching overhead. If you replace min/max expression `w = min(w, W-1);` with `if (w >= W) return;`, runtime increases up to 2x against to code without branching (initial version).<br>
|
||||
If branching is inevitable for your element-based kernel, it is recommended to change the scheme to line-based. See the kernel variant below:
|
||||
```cpp
|
||||
// Line-wise version
|
||||
__kernel void reorg(const __global half* restrict src, __global half* restrict out, int H, int W, int stride)
|
||||
{
|
||||
int h = min((int)get_global_id(0), H-1);
|
||||
|
||||
int c = get_global_id(1);
|
||||
int C = get_global_size(1);
|
||||
int C2 = C/(stride*stride);
|
||||
int offset = c / C2;
|
||||
int c2 = c - C2 * offset;
|
||||
|
||||
int H2 = H*stride;
|
||||
int W2 = W*stride;
|
||||
|
||||
for (int w = 0; w < W; ++w)
|
||||
{
|
||||
int h2 = h*stride + offset / stride;
|
||||
int w2 = w*stride + offset - stride * (offset / stride);
|
||||
|
||||
out[W*H*c + W*h + w] = src[W2*H2*c2 + W2*h2 + w2];
|
||||
}
|
||||
}
|
||||
```
|
||||
This decreases the execution time up to 40% against the best performing vectorized kernel without early exits (initial version).
|
||||
7. Reuse computations among work items by using line-based kernels or sharing values though `__local` memory.
|
||||
8. Improve data access locality. Most of custom kernels are memory bound while convolution and fully connected layers are hardware-implemented. The code below demonstrates a further optimized version of the `reorg` kernel unrolled by `stride`:
|
||||
```cpp
|
||||
// Unrolled line-wise version
|
||||
__kernel void reorg_unrolled_by_stride(const __global half* restrict src, __global half* restrict dst,
|
||||
int H, int W, int stride)
|
||||
{
|
||||
int h = min((int)get_global_id(0), H-1);
|
||||
|
||||
int c2 = get_global_id(1);
|
||||
int C2 = get_global_size(1);
|
||||
int C = C2*stride*stride;
|
||||
|
||||
int H2 = H*stride;
|
||||
int W2 = W*stride;
|
||||
|
||||
for (int stride_y = 0; stride_y < stride; stride_y++)
|
||||
for (int stride_x = 0; stride_x < stride; stride_x++)
|
||||
for (int w2 = 0, w = 0; w < W; w2 += stride, w++)
|
||||
dst[W*H*C2*(stride_y*stride+stride_x) + W*H*c2 + W*h + w] = src[W2*H2*c2 + W2*h*stride + W2*stride_y + w2 + stride_x];
|
||||
}
|
||||
```
|
||||
`scr` data in this case loaded only once. As the result, the cycle count drops up to 45% against the line-wise version.
|
||||
|
||||
9. Copy data from `__dlobal` to `__local` or `__private` memory if the data is accessed more than once. Access to
|
||||
`__dlobal` memory is orders of magnitude slower than access to `__local`/`__private` due to statically scheduled pipeline, which
|
||||
stalls completely on memory access without any prefetch. The same recommendation is applicable for scalar load/store
|
||||
from/to a `__blobal` pointer since work-group copying could be done in a vector fashion.
|
||||
|
||||
10. Use a manual DMA extension. Local (on-chip) memory throughput is up to 24x higher than DDR throughput. Starting from OpenVINO™ 2020.1, VPU OpenCL features manual-DMA kernel extension to copy sub-tensor used by work group into local memory and performing compute without DDR evolved. Here is the simple GRN kernel implementation that runs over DDR. Local size is in the form (width of the input tensor, 1, 1) to define a large enough work group to get code automatically vectorized and unrolled, while global size is (width of the input tensor, height of the input tensor, 1):
|
||||
```cpp
|
||||
__kernel void grn_NCHW(
|
||||
__global const half* restrict src_data,
|
||||
__global half* restrict dst_data,
|
||||
int C,
|
||||
float bias)
|
||||
{
|
||||
float variance = bias + 1e-9f;
|
||||
|
||||
#pragma unroll 4
|
||||
for (int c = 0; c < C; c++)
|
||||
{
|
||||
float val = (float) src_data[c*get_global_size(1)*get_global_size(0) + get_global_id(1)*get_global_size(0) + get_global_id(0)];
|
||||
variance += val*val;
|
||||
}
|
||||
|
||||
half hvariance = (half)(native_rsqrt((half)(variance/16.f))*0.25f);
|
||||
|
||||
#pragma unroll 4
|
||||
for (int c = 0; c < C; c++)
|
||||
{
|
||||
dst_data[c*get_global_size(1)*get_global_size(0) + get_global_id(1)*get_global_size(0) + get_global_id(0)]
|
||||
= src_data[c*get_global_size(1)*get_global_size(0) + get_global_id(1)*get_global_size(0) + get_global_id(0)] * hvariance;
|
||||
}
|
||||
}
|
||||
```
|
||||
|
||||
This kernel can be rewritten to introduce special data binding `__dma_preload` and `__dma_postwrite intrinsics`. This means that instead of one kernel, a group of three kernels should be implemented: `kernelName`, `__dma_preload_kernelName`, and `__dma_postwrite_kernelName`. `__dma_preload_kernelName` for a particular work group `n` is guaranteed to be executed before the `n`-th work group itself, while `__dma_postwrite_kernelName` is guaranteed to be executed after a corresponding work group. You can define one of those functions that are intended to be used to copy data from-to `__global` and `__local` memory. The syntactics requires exact functional signature match. The example below illustrates how to prepare your kernel for manual-DMA.
|
||||
|
||||
```cpp
|
||||
__kernel void __dma_preload_grn_NCHW(
|
||||
__global const half* restrict src,
|
||||
__global half* restrict dst,
|
||||
__local half* restrict local_src,
|
||||
__local half* restrict local_dst,
|
||||
int C,
|
||||
float bias)
|
||||
{
|
||||
// ToDO: copy required piece of src tensor into local_src
|
||||
}
|
||||
|
||||
__kernel void __dma_postwrite_grn_NCHW(
|
||||
__global const half* restrict src,
|
||||
__global half* restrict dst,
|
||||
__local const half* restrict local_src,
|
||||
__local half* restrict local_dst,
|
||||
int C,
|
||||
float bias)
|
||||
{
|
||||
// ToDO: copy back computed piece of local_dst into dst
|
||||
}
|
||||
|
||||
__kernel void grn_NCHW(
|
||||
__global const half* restrict src_data,
|
||||
__global half* restrict dst_data,
|
||||
__local half* restrict src,
|
||||
__local half* restrict dst,
|
||||
int C,
|
||||
float bias)
|
||||
{
|
||||
// same as the example above
|
||||
}
|
||||
```
|
||||
The GRN kernel operates on channel-major tensors to compute average over full channel range and then normalizes input elements to produce the output.
|
||||
As a part of the manual DMA extension, a group of work group copy functions are introduced in addition to `async_work_group_copy`, which is also mapped to a DMA call.
|
||||
|
||||
Here is the list of supported functions:
|
||||
```cpp
|
||||
// 2D sub-tensor copy
|
||||
event_t WorkGroupDmaCreateStrideTransaction(
|
||||
const local T *src,
|
||||
global T *dst,
|
||||
size_t src_width, // width of the line of source in bytes
|
||||
size_t dst_width, // width of the line of destination in bytes
|
||||
size_t src_stride, // stride between corresponding 2 consecutive lines of source in bytes
|
||||
size_t dst_stride, // stride between corresponding 2 consecutive lines of destination in bytes
|
||||
size_t size, // total number of bytes loaded for all lines from source to destination
|
||||
event_t event) __OVERLOAD;
|
||||
|
||||
|
||||
event_t WorkGroupDmaCreateStrideTransaction(
|
||||
const global T *src,
|
||||
local T *dst,
|
||||
size_t src_width, // width of the line of source in bytes
|
||||
size_t dst_width, // width of the line of destination in bytes
|
||||
size_t src_stride, // stride between corresponding 2 consecutive lines of source in bytes
|
||||
size_t dst_stride, // stride between corresponding 2 consecutive lines of destination in bytes
|
||||
size_t size, // total number of bytes loaded for all lines from source to destination
|
||||
event_t event) __OVERLOAD;
|
||||
|
||||
// 3D sub-tensor copy
|
||||
event_t WorkGroupDmaCreate3DTransaction(
|
||||
const local T *src,
|
||||
global T *dst,
|
||||
size_t src_width, // width of the line of source in bytes
|
||||
size_t dst_width, // width of the line of destination in bytes
|
||||
size_t src_stride, // stride between corresponding 2 consecutive lines of source in bytes
|
||||
size_t dst_stride, // stride between corresponding 2 consecutive lines of destination in bytes
|
||||
size_t num_planes, // number of planes to be copied
|
||||
size_t src_plane_stride, // stride between corresponding 2 consecutive planes of source in bytes
|
||||
size_t dst_plane_stride, // stride between corresponding 2 consecutive planes of destination in bytes
|
||||
size_t size, // size of the loaded plane in bytes, analogues to the size in 2D case
|
||||
event_t event) __OVERLOAD;
|
||||
|
||||
event_t WorkGroupDmaCreate3DTransaction(
|
||||
const global T *src,
|
||||
local T *dst,
|
||||
size_t src_width, // width of the line of source in bytes
|
||||
size_t dst_width, // width of the line of destination in bytes
|
||||
size_t src_stride, // stride between corresponding 2 consecutive lines of source in bytes
|
||||
size_t dst_stride, // stride between corresponding 2 consecutive lines of destination in bytes
|
||||
size_t num_planes, // number of planes to be copied
|
||||
size_t src_plane_stride, // stride between corresponding 2 consecutive planes of source in bytes
|
||||
size_t dst_plane_stride, // stride between corresponding 2 consecutive planes of destination in bytes
|
||||
size_t size, // size of the loaded plane in bytes, analogues to the size in 2D case
|
||||
event_t event) __OVERLOAD;
|
||||
```
|
||||
where `T` can be `uchar`, `char`, `short`, `ushort`, `int`, `uint`, `long`, `ulong`, `half` or `float`.
|
||||
|
||||
Modified version of the GRN kernel could be the following:
|
||||
```cpp
|
||||
__kernel void __dma_preload_grn_NCHW(
|
||||
__global const half* restrict src,
|
||||
__global half* restrict dst,
|
||||
__local half* restrict local_src,
|
||||
__local half* restrict local_dst,
|
||||
int C,
|
||||
float bias)
|
||||
{
|
||||
WorkGroupDmaCreate3DTransaction(
|
||||
src + get_group_id(0)*get_local_size(0)
|
||||
+ get_group_id(1)*get_local_size(1)*get_global_size(0), // src
|
||||
local_src, // dst
|
||||
get_local_size(0) * sizeof(half), // src width
|
||||
get_local_size(0) * sizeof(half), // dst width
|
||||
get_global_size(0) * sizeof(half), // src stride
|
||||
get_local_size(0) * sizeof(half), // dst stride
|
||||
C, // num planes
|
||||
get_global_size(0) * get_global_size(1) * sizeof(half), // src plane stride
|
||||
get_local_size(0) * get_local_size(1) * sizeof(half), // dst plane stride
|
||||
get_local_size(0) * get_local_size(1) * sizeof(half), // plane size
|
||||
0);
|
||||
}
|
||||
|
||||
__kernel void __dma_postwrite_grn_NCHW(
|
||||
__global const half* restrict src,
|
||||
__global half* restrict dst,
|
||||
__local const half* restrict local_src,
|
||||
__local half* restrict local_dst,
|
||||
int C,
|
||||
float bias)
|
||||
{
|
||||
WorkGroupDmaCreate3DTransaction(
|
||||
local_dst, // src
|
||||
dst + get_group_id(0)*get_local_size(0)
|
||||
+ get_group_id(1)*get_local_size(1)*get_global_size(0), // dst
|
||||
get_local_size(0) * sizeof(half), // src width
|
||||
get_local_size(0) * sizeof(half), // dst width
|
||||
get_local_size(0) * sizeof(half), // src stride
|
||||
get_global_size(0) * sizeof(half), // dst stride
|
||||
C, // num planes
|
||||
get_local_size(0) * get_local_size(1) * sizeof(half), // src plane stride
|
||||
get_global_size(0) * get_global_size(1) * sizeof(half), // dst plane stride
|
||||
get_local_size(0) * get_local_size(1) * sizeof(half), // plane size
|
||||
0);
|
||||
}
|
||||
|
||||
__kernel void grn_NCHW(
|
||||
__global const half* restrict src_data,
|
||||
__global half* restrict dst_data,
|
||||
__local half* restrict src,
|
||||
__local half* restrict dst,
|
||||
int C,
|
||||
float bias)
|
||||
{
|
||||
float variance = bias + 1e-9f;
|
||||
|
||||
#pragma unroll 8
|
||||
for (int c = 0; c < C; c++)
|
||||
{
|
||||
float val = (float) src[c*get_local_size(1)*get_local_size(0) + get_local_id(1)*get_local_size(0) + get_local_id(0)];
|
||||
variance += val*val;
|
||||
}
|
||||
|
||||
half hvariance = (half)(native_rsqrt((half)(variance/16.f))*0.25f);
|
||||
|
||||
#pragma unroll 8
|
||||
for (int c = 0; c < C; c++)
|
||||
{
|
||||
dst[c*get_local_size(1)*get_local_size(0) + get_local_id(1)*get_local_size(0) + get_local_id(0)]
|
||||
= src[c*get_local_size(1)*get_local_size(0) + get_local_id(1)*get_local_size(0) + get_local_id(0)] * hvariance;
|
||||
}
|
||||
}
|
||||
```
|
||||
|
||||
Note the `get_local_size` and `get_local_id` usage inside the kernel. 21x speedup is expected for a kernel on enet-curbs setup because it was completely limited by memory usage.
|
||||
|
||||
An alternative method to using DMA is to use work item copy extension. Those functions are executed inside a kernel and requires work groups equal to single work item.
|
||||
|
||||
Here is the list of supported work item functions:
|
||||
```cpp
|
||||
item_dma_event_t WorkItemDmaCreateTransaction(
|
||||
const global T *src,
|
||||
private T *dst,
|
||||
size_t size,
|
||||
item_dma_event_t event) __OVERLOAD;
|
||||
|
||||
item_dma_event_t WorkItemDmaCreateTransaction(
|
||||
const private T *src,
|
||||
global T *dst,
|
||||
size_t size,
|
||||
item_dma_event_t event) __OVERLOAD;
|
||||
|
||||
item_dma_event_t WorkItemDmaCreateStrideTransaction(
|
||||
const global T *src,
|
||||
private T *dst,
|
||||
size_t src_width,
|
||||
size_t dst_width,
|
||||
size_t src_stride,
|
||||
size_t dst_stride,
|
||||
size_t size,
|
||||
item_dma_event_t event) __OVERLOAD;
|
||||
|
||||
item_dma_event_t WorkItemDmaCreateStrideTransaction(
|
||||
const private T *src,
|
||||
global T *dst,
|
||||
size_t src_width,
|
||||
size_t dst_width,
|
||||
size_t src_stride,
|
||||
size_t dst_stride,
|
||||
size_t size,
|
||||
item_dma_event_t event) __OVERLOAD;
|
||||
|
||||
item_dma_event_t WorkItemDmaCreate3DTransaction(
|
||||
const global T *src,
|
||||
private T *dst,
|
||||
size_t src_width,
|
||||
size_t dst_width,
|
||||
size_t src_stride,
|
||||
size_t dst_stride,
|
||||
size_t num_planes,
|
||||
size_t src_plane_stride,
|
||||
size_t dst_plane_stride,
|
||||
size_t size,
|
||||
item_dma_event_t event) __OVERLOAD;
|
||||
|
||||
item_dma_event_t WorkItemDmaCreate3DTransaction(
|
||||
const private T *src,
|
||||
global T *dst,
|
||||
size_t src_width,
|
||||
size_t dst_width,
|
||||
size_t src_stride,
|
||||
size_t dst_stride,
|
||||
size_t num_planes,
|
||||
size_t src_plane_stride,
|
||||
size_t dst_plane_stride,
|
||||
size_t size,
|
||||
item_dma_event_t event) __OVERLOAD;
|
||||
```
|
||||
where `T` can be `uchar`, `char`, `short`, `ushort`, `int`, `uint`, `long`, `ulong`, `half` or `float`.
|
||||
@@ -112,7 +112,7 @@ To keep the model valid after the reshape, choose a new input shape that satisfi
|
||||
For details, refer to the <a href="_docs_MO_DG_prepare_model_convert_model_tf_specific_Convert_Object_Detection_API_Models.html#tf_od_custom_input_shape">Tensorflow Object Detection API models resizing techniques</a>.
|
||||
|
||||
### Extensibility
|
||||
The Inference Engine provides a special mechanism that allows adding support of shape inference for custom operations. This mechanism is described in the [Extensibility documentation](Extensibility_DG/Intro.md)
|
||||
The Inference Engine provides a special mechanism that allows adding support of shape inference for custom operations. This mechanism is described in the [Extensibility documentation](../Extensibility_UG/Intro.md)
|
||||
|
||||
## Introduction (Python)
|
||||
|
||||
@@ -218,7 +218,7 @@ exec_net = ie.load_network(network=net, device_name="CPU")
|
||||
```
|
||||
|
||||
### Extensibility
|
||||
The Inference Engine provides a special mechanism that allows adding support of shape inference for custom operations. This mechanism is described in the [Extensibility documentation](Extensibility_DG/Intro.md)
|
||||
The Inference Engine provides a special mechanism that allows adding support of shape inference for custom operations. This mechanism is described in the [Extensibility documentation](../Extensibility_UG/Intro.md)
|
||||
|
||||
### See Also:
|
||||
|
||||
|
||||
@@ -14,7 +14,7 @@ For details on how to build a model in OpenVINO™ Runtime, see the [Build a Mod
|
||||
|
||||
## Operations
|
||||
|
||||
The `ov::Op` class represents any abstract operation in the model representation. Use this class to create [custom operations](../OV_Runtime_UG/Extensibility_DG/AddingNGraphOps.md).
|
||||
The `ov::Op` class represents any abstract operation in the model representation. Use this class to create [custom operations](../Extensibility_UG/add_openvino_ops).
|
||||
|
||||
## Operation Sets
|
||||
|
||||
@@ -39,7 +39,7 @@ Operation set `opsetX` integrates a list of pre-compiled operations that work
|
||||
|
||||
For a complete list of operation sets supported in OpenVINO™ toolkit, see [Available Operations Sets](../ops/opset.md).
|
||||
|
||||
To add suport of custom operations, see the [Add Custom OpenVINO Operations](../OV_Runtime_UG/Extensibility_DG/Intro.md) document.
|
||||
To add support of custom operations, see the [Add Custom OpenVINO Operations](../Extensibility_UG/Intro.md) document.
|
||||
|
||||
To build an `ov::Model` instance from `opset8` operations, include the following files:
|
||||
|
||||
@@ -88,4 +88,4 @@ The following code creates a model with several outputs:
|
||||
## See Also
|
||||
|
||||
* [Available Operation Sets](../ops/opset.md)
|
||||
* [OpenVINO™ Runtime Extensibility Developer Guide](../OV_Runtime_UG/Extensibility_DG/Intro.md)
|
||||
* [OpenVINO™ Runtime Extensibility Developer Guide](../Extensibility_UG/Intro.md)
|
||||
@@ -13,7 +13,6 @@
|
||||
openvino_docs_IE_DG_Int8Inference
|
||||
openvino_docs_IE_DG_Bfloat16Inference
|
||||
openvino_docs_transformations
|
||||
openvino_docs_IE_DG_Extensibility_DG_Intro
|
||||
|
||||
@endsphinxdirective
|
||||
|
||||
|
||||
@@ -6,7 +6,7 @@ The CPU plugin was developed to achieve high performance of neural networks on C
|
||||
|
||||
Currently, the CPU plugin uses Intel® Threading Building Blocks (Intel® TBB) in order to parallelize calculations. Please refer to the [Optimization Guide](../../optimization_guide/dldt_optimization_guide.md) for associated performance considerations.
|
||||
|
||||
The set of supported layers can be expanded with [the Extensibility mechanism](../Extensibility_DG/Intro.md).
|
||||
The set of supported layers can be expanded with [the Extensibility mechanism](../../Extensibility_UG/Intro.md).
|
||||
|
||||
## Supported Platforms
|
||||
|
||||
|
||||
@@ -271,6 +271,6 @@ The following layers are supported by the plugins and by [Shape Inference featur
|
||||
|
||||
\*- support is limited to the specific parameters. Refer to "Known Layers Limitation" section for the device [from the list of supported](Supported_Devices.md).
|
||||
|
||||
\*\*- support is implemented via [Extensibility mechanism](../Extensibility_DG/Intro.md).
|
||||
\*\*- support is implemented via [Extensibility mechanism](../../Extensibility_UG/Intro.md).
|
||||
|
||||
\*\*\*- supports NCDHW layout.
|
||||
@@ -138,7 +138,7 @@ In a perfect pipeline this time should be near zero, which means that the data w
|
||||
|
||||
**Get the following message when running inference with the VPU plugin: "[VPU] Cannot convert layer <layer_name> due to unsupported layer type <layer_type>"**
|
||||
|
||||
This means that your topology has a layer that is unsupported by your target VPU plugin. To resolve this issue, you can implement the custom layer for the target device using the [Inference Engine Extensibility mechanism](../Extensibility_DG/Intro.md). Or, to quickly get a working prototype, you can use the heterogeneous scenario with the default fallback policy (see the [Heterogeneous execution](../hetero_execution.md) section). Use the HETERO mode with a fallback device that supports this layer, for example, CPU: `HETERO:MYRIAD,CPU`.
|
||||
This means that your topology has a layer that is unsupported by your target VPU plugin. To resolve this issue, you can implement the custom layer for the target device using the [OpenVINO™ Extensibility mechanism](../../Extensibility_UG/Intro.md). Or, to quickly get a working prototype, you can use the heterogeneous scenario with the default fallback policy (see the [Heterogeneous execution](../hetero_execution.md) section). Use the HETERO mode with a fallback device that supports this layer, for example, CPU: `HETERO:MYRIAD,CPU`.
|
||||
For a list of VPU-supported layers, see the Supported Layers section of the [Supported Devices](Supported_Devices.md) page.
|
||||
|
||||
## Known Layers Limitations
|
||||
|
||||
@@ -8,7 +8,6 @@
|
||||
:hidden:
|
||||
|
||||
openvino_docs_MO_DG_Deep_Learning_Model_Optimizer_DevGuide
|
||||
openvino_docs_HOWTO_Custom_Layers_Guide
|
||||
omz_tools_downloader
|
||||
|
||||
|
||||
@@ -68,9 +67,10 @@
|
||||
|
||||
.. toctree::
|
||||
:maxdepth: 1
|
||||
:caption: Developing Inference Engine Plugins
|
||||
:caption: OpenVINO Extensibility
|
||||
:hidden:
|
||||
|
||||
openvino_docs_Extensibility_UG_Intro
|
||||
Inference Engine Plugin Developer Guide <openvino_docs_ie_plugin_dg_overview>
|
||||
groupie_dev_api
|
||||
Plugin Transformation Pipeline <openvino_docs_IE_DG_plugin_transformation_pipeline>
|
||||
|
||||
@@ -13,7 +13,7 @@
|
||||
* [Accelerate Vision-based AI with Intel® Distribution of OpenVINO™ Toolkit](https://www.intel.ai/accelerate-vision-based-ai-with-intel-distribution-of-openvino-toolkit/)
|
||||
|
||||
## Custom Operations Guide
|
||||
To learn about what is *custom operation* and how to work with them in the Deep Learning Deployment Toolkit, see the [Custom Operations Guide](../HOWTO/Custom_Layers_Guide.md).
|
||||
To learn about what is *custom operation* and how to work with them in the Deep Learning Deployment Toolkit, see the [Custom Operations Guide](../Extensibility_UG/Intro.md).
|
||||
|
||||
## Introducing OpenVINO™ and Computer Vision | IoT Developer Show Season 2 | Intel Software
|
||||
|
||||
|
||||
@@ -35,6 +35,7 @@ list(REMOVE_ITEM SOURCES "${CMAKE_CURRENT_SOURCE_DIR}/dldt_optimization_guide2.c
|
||||
# create a static library
|
||||
|
||||
add_library(${TARGET_NAME} STATIC ${SOURCES})
|
||||
target_include_directories(${TARGET_NAME} PRIVATE "${CMAKE_CURRENT_SOURCE_DIR}/../template_extension/new/")
|
||||
|
||||
if(CLDNN__IOCL_ICD_INCDIRS)
|
||||
target_include_directories(${TARGET_NAME} SYSTEM PRIVATE ${CLDNN__IOCL_ICD_INCDIRS})
|
||||
|
||||
@@ -0,0 +1,26 @@
|
||||
// Copyright (C) 2018-2021 Intel Corporation
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
//
|
||||
#include <openvino/openvino.hpp>
|
||||
#include <openvino/core/op_extension.hpp>
|
||||
#include <identity.hpp>
|
||||
|
||||
int main() {
|
||||
{
|
||||
//! [add_extension]
|
||||
ov::Core core;
|
||||
// Use operation type to add operation extension
|
||||
core.add_extension<TemplateExtension::Identity>();
|
||||
// or you can add operation extension to this method
|
||||
core.add_extension(ov::OpExtension<TemplateExtension::Identity>());
|
||||
//! [add_extension]
|
||||
}
|
||||
{
|
||||
//! [add_extension_lib]
|
||||
ov::Core core;
|
||||
// Load extensions library to ov::Core
|
||||
core.add_extension("openvino_template_extension.so");
|
||||
//! [add_extension_lib]
|
||||
}
|
||||
return 0;
|
||||
}
|
||||
@@ -0,0 +1,15 @@
|
||||
# Copyright (C) 2018-2022 Intel Corporation
|
||||
# SPDX-License-Identifier: Apache-2.0
|
||||
#
|
||||
|
||||
import openvino.runtime as ov
|
||||
|
||||
#! [add_extension]
|
||||
# Not implemented
|
||||
#! [add_extension]
|
||||
|
||||
#! [add_extension_lib]
|
||||
core = ov.Core()
|
||||
# Load extensions library to ov::Core
|
||||
core.add_extension("openvino_template_extension.so")
|
||||
#! [add_extension_lib]
|
||||
@@ -14,7 +14,13 @@ set(SRC identity.cpp ov_extension.cpp)
|
||||
add_library(${TARGET_NAME} MODULE ${SRC})
|
||||
|
||||
target_compile_definitions(${TARGET_NAME} PRIVATE IMPLEMENT_OPENVINO_EXTENSION_API)
|
||||
target_link_libraries(${TARGET_NAME} PRIVATE openvino::core)
|
||||
target_link_libraries(${TARGET_NAME} PRIVATE openvino::runtime)
|
||||
|
||||
# To map custom operation to framework
|
||||
if(OpenVINO_Frontend_ONNX_FOUND)
|
||||
target_link_libraries(${TARGET_NAME} PRIVATE openvino::frontend::onnx)
|
||||
target_compile_definitions(${TARGET_NAME} PRIVATE OPENVINO_ONNX_FRONTEND_ENABLED)
|
||||
endif()
|
||||
# [cmake:extension]
|
||||
|
||||
# Enable code style check
|
||||
|
||||
@@ -4,7 +4,14 @@
|
||||
|
||||
#pragma once
|
||||
|
||||
//! [op:common_include]
|
||||
#include <openvino/op/op.hpp>
|
||||
//! [op:common_include]
|
||||
//! [op:frontend_include]
|
||||
#ifdef OPENVINO_ONNX_FRONTEND_ENABLED
|
||||
# include <openvino/frontend/onnx/extension/op.hpp>
|
||||
#endif
|
||||
//! [op:frontend_include]
|
||||
|
||||
//! [op:header]
|
||||
namespace TemplateExtension {
|
||||
@@ -13,6 +20,10 @@ class Identity : public ov::op::Op {
|
||||
public:
|
||||
OPENVINO_OP("Identity");
|
||||
|
||||
#ifdef OPENVINO_ONNX_FRONTEND_ENABLED
|
||||
OPENVINO_FRAMEWORK_MAP(onnx)
|
||||
#endif
|
||||
|
||||
Identity() = default;
|
||||
Identity(const ov::Output<ov::Node>& arg);
|
||||
void validate_and_infer_types() override;
|
||||
|
||||
@@ -7,5 +7,11 @@
|
||||
|
||||
#include "identity.hpp"
|
||||
|
||||
// clang-format off
|
||||
//! [ov_extension:entry_point]
|
||||
OPENVINO_CREATE_EXTENSIONS(
|
||||
std::vector<ov::Extension::Ptr>({std::make_shared<ov::OpExtension<TemplateExtension::Identity>>()}));
|
||||
std::vector<ov::Extension::Ptr>({
|
||||
std::make_shared<ov::OpExtension<TemplateExtension::Identity>>()
|
||||
}));
|
||||
//! [ov_extension:entry_point]
|
||||
// clang-format on
|
||||
@@ -99,7 +99,7 @@ public:
|
||||
const std::string& target_name = p_name != m_attr_names_map.end() ? p_name->second : name;
|
||||
try {
|
||||
adapter.set_as_any(m_context.get_attribute_as_any(target_name));
|
||||
} catch (::ov::AssertFailure ex) {
|
||||
} catch (::ov::AssertFailure& ex) {
|
||||
OPENVINO_ASSERT(false,
|
||||
ex.what(),
|
||||
"\nValue for attribute \"",
|
||||
@@ -259,4 +259,4 @@ using OpExtension = ov::frontend::OpExtensionBase<ov::frontend::ConversionExtens
|
||||
};
|
||||
|
||||
} // namespace frontend
|
||||
} // namespace ov
|
||||
} // namespace ov
|
||||
Reference in new issue
Block a user