How to Implement Custom GPU Operations — OpenVINO™ documentation (original) (raw)

To enable operations not supported by OpenVINO™ out of the box, you may need an extension for OpenVINO operation set, and a custom kernel for the device you will target. This article 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 automatically-loaded<lib_path>/cldnn_global_custom_kernels/cldnn_global_custom_kernels.xml file.
• Call the ov::Core::set_property() method from your application with the"CONFIG_FILE" key and the configuration file name as a value before loading the network that uses custom operations to the plugin:

Python

core = ov.Core()
core.set_property("GPU", {"CONFIG_FILE": "<path_to_the_xml_file>"})

C++

ov::Core core;
// Load GPU Extensions
core.set_property("GPU", {{ "CONFIG_FILE", "<path_to_the_xml_file>" }});

All OpenVINO 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:

$ ./classification_sample -m /bvlc_alexnet_fp16.xml -i ./validation_set/daily/227x227/apron.bmp -d GPU -c /custom_layer_example.xml

Important

Due to the deprecation of Open Model Zoo, models in the OpenVINO IR format are now published on Hugging Face.

Configuration File Format#

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:

CustomLayer Node and Sub-Node Structure#

The CustomLayer node contains the entire configuration for a single custom operation.

Sub-nodes: Kernel (1), Buffers (1), CompilerOptions (0+),WorkSizes (0/1)

Kernel Node and Sub-Node Structure#

The Kernel node contains all kernel source code configuration.

Sub-nodes: Source (1+), Define (0+)

Source Node and Sub-Node Structure#

The Source node points to a single OpenCL source file.

Sub-nodes: None

Define Node and Sub-Node Structure#

The Define node configures a single #define instruction to be added to the sources during compilation (JIT).

Sub-nodes: None

The resulting JIT has the following form:#define [name] [type] [value/default].

Buffers Node and Sub-Node Structure#

The 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#

The Data node configures a single input with static data, for example, weights or biases.

Sub-nodes: None

Tensor Node and Sub-Node Structure#

The Tensor node configures a single input or output tensor.

CompilerOptions Node and Sub-Node Structure#

The CompilerOptions node configures the compilation flags for the OpenCL sources.

Sub-nodes: None

WorkSizes Node and Sub-Node Structure#

The WorkSizes node configures the global/local work sizes to be used when queuing an OpenCL program for execution.

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 theConfiguration File Format.

Built-In Definitions for Custom Layers#

The following table includes definitions that are attached before user sources.

For an example, see Example Kernel.

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:

#define INPUT0_DIMS_SIZE 4 #define INPUT0_DIMS (int []){ 1,96,55,55, }

Example Kernel#

#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); // batchesfeatures, 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 = batchINPUT0_PITCHES[0] + featureINPUT0_PITCHES[1] + idyINPUT0_PITCHES[2] + idxINPUT0_PITCHES[3] + INPUT0_OFFSET; const uint out_id = batchOUTPUT0_PITCHES[0] + featureOUTPUT0_PITCHES[1] + idyOUTPUT0_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 such as the INPUT0_TYPE are actually defined as OpenCL (pre-)compiler inputs by OpenVINO for efficiency reasons. See the Debugging Tips below for information on debugging the results.

Debugging Tips#

Using printf in the OpenCL™ Kernels.

To debug the specific values, 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.

For more information, refer to theprintf Function.