OpenVX Vision Image Extension API Introduction - Basic API

Document created by Devin Jiao Employee on Feb 18, 2020Last modified by Devin Jiao Employee on Feb 24, 2020
Version 3Show Document
  • View in full screen mode

Introduction

EVIS (Enhanced Vision Instruction Set) is an API level program language, which is applicable on GC7000XSVX (i.MX8QM) and VIP8000NanoSi (i.MX8MP). The instructions take advantage of the enhanced vision capabilities in the vision-capble hardware, with low-latency. It provides additional functionality for vision image processing beyond the functions provided through the Khronos Group OpenVX API. In the i.MX8QM/i.MX8MP BSP, the OpenCL compiler also support the EVIS instructions. So, it is also an OpenCL VX Extension API.

The source file can found in /usr/include/CL/cl_viv_vx_ext.h.

 

Hardware Requirements

i.MX8QM (GC7000XSVX) supports EVIS1.

i.MX8MP (VIP8000NanoSi) supports EVIS2.

 

Packed Data Types

To fully utilize the computing power of the instructions, the extension API support packed data types. For example, in standard OpenCL, a vector char4 will occupy 4x 32-bit registers, while a packed char16 only occupies 128 bits. Thus use packed data types as possible. 

The char, unsigned char, short, unsigned short, integer, unsigned integer, float packed data types are supported. They are defined with vxc_ prefix i.e. vxc_char, vxc_uchar, vxc_short, vxc_ushort, vxc_int, vxc_uint, vxc_float, followed by a literal value n that defines the number of elements in the packed data. Supported values of n are 2, 4, 8, and 16 for all the packed data types.

Table 1 List of packed data type

TypeDescription
vxc_charnA vector of n packed signed character value
vxc_ucharnA vector of n packed unsigned character value
vxc_shortnA vector of n packed signed short value
vxc_ushortnA vector of n packed unsigned short value
vxc_intnA vector of n packed signed integer value
vxc_uintnA vector of n packed unsigned integer value
vxc_floatnA vector of n packed float value

 

OP_CODE Instructions

OP_CODE instructions operate on packed data. The enumeration can be found in /usr/include/CL/cl_viv_vx_ext.h.

Only EVIS1 supports instructions:

VXC_IAdd

VXC_MagPhase

VXC_BiLinear

VXC_SelectAdd

VXC_BitReplace

VXC_Filter

VXC_DP2x16/VXC_DP2x16_b

 

  • Objects load and store

Packed type image data read/write: supported types are packed 8-bit/16bit integer, 16bit float. Image read/write for image1d_t/image1d_array/image2d_t. Offset should be composed by using VXC_5BITOFFSET_XY(x, y).

VXC_OP4(img_load, Dest, Image, Coord, Offset, Info)

VXC_OP4_NoDest(img_store, Image, Coord, Color, Info)

Parameters:

        img_load/img_store    Read/write image data. 

        Dest                            The destination loading the data to.

        Image                          The packed image data read from for img_load. The packed image data writing to for img_store.

        Coord                          Coordinates to read/write the image data.

        Color                           The image data being written to Image for img_store.

        Info                              See more info in VXC_MODIFIER(StartBin, EndBin, SourceBin, RoundingMode, Clamp).

VXC_MODIFIER(StartBin, EndBin, SourceBin, RoundingMode, Clamp)

Parameters:

        StartBin/EndBin           The first bin/the last bin for consecutive packed data.

        SourceBin                    Not used. 

        RoundingMode            0: Toward Zero (truncated), 1: Toward Infinity (rounded up), 2: To Nearest Even, 3: not used.

        Clamp                          0: no, result is truncated to fit result type (just the lower bits are copied), 1: yes, result is clamped to fit the result type.

For example,

int2 coord = (int2)(get_global_id(0), get_global_id(1));

vxc_uchar16 r1;

VXC_OP4(img_load, r1, in_image, coord, 0, VXC_MODIFIER(0, 15, 0, VXC_RM_TowardZero, 0));

 

  • VXC_Filter

This interface applies a specified filter on a 3x3 pixel block.

VXC_OP4(filter, Dest, Src0, Src1, Src2, Info)

Parameters:

        filter                           Filter modes.

        Dest                          The filtered image.

        Src0                          The first row pixels for 3x3 filter.

        Src1                          The second row pixels for 3x3 filter.

        Src2                          The third row pixles for 3x3 filter.

        Info                            See more info in VXC_MODIFIER_FILTER(StartBin, EndBin, SourceBin, Filter, Clamp).

VXC_MODIFIER_FILTER(StartBin, EndBin, SourceBin, Filter, Clamp)

Parameters:

        StartBin/EndBin        The first bin/the last bin for consecutive packed data.

        SourceBin                 Not used.

        Filter                          Filter modes are listed in table 2.

        Clamp                        0: no, result is truncated to fit result type (just the lower bits are copied), 1: yes, result is clamped to fit the result type.

 

Table 2. List of filter modes:

Filter ModeDescription
VXC_FM_BOXCompute a 3x3 box filter:
|1/9, 1/9, 1/9, 1/9, 1/9, 1/9, 1/9, 1/9, 1/9|.
VXC_FM_GuassianCompute a 3x3 Gaussian filter:
|1/16, 2/16, 1/16, 2/16, 4/16, 2/16, 1/16, 2/16, 1/16|.
VXC_FM_SobelXCompute a 3x3 Sobel filter in the x-direction:
|-1, 0, 1, -2, 0, 2, -1, 0, 1|.
VXC_FM_SobelYCompute a 3x3 Sobel filter in the y-direction:
|-1, -2, -1, 0, 0, 0, 1, 2, 1|.
VXC_FM_ScharrXCompute a 3x3 Scharr filter in the x-direction:
|3, 0, -3, 10, 0, -10, 3, 0, -3|.
VXC_FM_ScharrYCompute a 3x3 Scharr filter in the y-direction:
|3, 10, 3, 0, 0, 0, -3, -10, -3|.
VXC_FM_MaxGet the maximum from a 3x3 kernel.
VXC_FM_MinGet the minimum from a 3x3 kernel.
VXC_FM_MedianGet the median from a 3x3 kernel.

 

For example (details in Gaussian Filter examples),

int2 coord_in1 = coord + (int2)(-1, -1);\n\
VXC_OP4(img_load, lineA, in_image, coord_in1, 0, VXC_MODIFIER(0, 15, 0, VXC_RM_TowardZero, 0));\n\
int2 coord_in2 = coord + (int2)(-1, 0);\n\
VXC_OP4(img_load, lineB, in_image, coord_in2, 0, VXC_MODIFIER(0, 15, 0, VXC_RM_TowardZero, 0));\n\
int2 coord_in3 = coord + (int2)(-1, 1);\n\
VXC_OP4(img_load, lineC, in_image, coord_in3, 0, VXC_MODIFIER(0, 15, 0, VXC_RM_TowardZero, 0));\n\
int info = VXC_MODIFIER_FILTER(0, 13, 0, VXC_FM_Guassian, 0);\n\
VXC_OP4(filter, out, lineA, lineB, lineC, info); ;\n\

  • VXC_AbsDiff

Calculates a result for the absolute difference between a and b. It works on packed data, so it can compute 16x 8-bit values or 8x 16-bit values.

VXC_OP3(abs_diff, Dest, Src0, Src1, Info)

Parameters:

        abs_diff                 Specify the function of absolute difference.

        Dest                      Destination to store the result.

        Src0                      The first source to calculate the absolute difference.

        Src1                      The second source to calculate the absolute differenece.

        Info                       See more info in VXC_MODIFIER(StartBin, EndBin, SourceBin, RoundingMode, Clamp).

 

There are also other interfaces will not be specified here, which can be found in the /usr/include/CL/cl_viv_vx_ext.h, i.e. VXC_IAdd, VXC_IAccSq, VXC_Lerp, VXC_MagPhase, VXC_MulShift, VXC_Clamp, VXC_BiLinear, VXC_SelectAdd, VXC_AtomicAdd, VXC_BitExtract and VXC_BitReplace. 

 

Further Reading:

OpenVX Vision Image Extension API Introduction - DP Dot Products

2 people found this helpful

Attachments

    Outcomes