OpenVX Vision Image Extension API Introduction - Basic API

5 Kudos

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

Type	Description
vxc_charn	A vector of n packed signed character value
vxc_ucharn	A vector of n packed unsigned character value
vxc_shortn	A vector of n packed signed short value
vxc_ushortn	A vector of n packed unsigned short value
vxc_intn	A vector of n packed signed integer value
vxc_uintn	A vector of n packed unsigned integer value
vxc_floatn	A 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 Mode	Description
VXC_FM_BOX	Compute a 3x3 box filter: \|1/9, 1/9, 1/9, 1/9, 1/9, 1/9, 1/9, 1/9, 1/9\|.
VXC_FM_Guassian	Compute a 3x3 Gaussian filter: \|1/16, 2/16, 1/16, 2/16, 4/16, 2/16, 1/16, 2/16, 1/16\|.
VXC_FM_SobelX	Compute a 3x3 Sobel filter in the x-direction: \|-1, 0, 1, -2, 0, 2, -1, 0, 1\|.
VXC_FM_SobelY	Compute a 3x3 Sobel filter in the y-direction: \|-1, -2, -1, 0, 0, 0, 1, 2, 1\|.
VXC_FM_ScharrX	Compute a 3x3 Scharr filter in the x-direction: \|3, 0, -3, 10, 0, -10, 3, 0, -3\|.
VXC_FM_ScharrY	Compute a 3x3 Scharr filter in the y-direction: \|3, 10, 3, 0, 0, 0, -3, -10, -3\|.
VXC_FM_Max	Get the maximum from a 3x3 kernel.
VXC_FM_Min	Get the minimum from a 3x3 kernel.
VXC_FM_Median	Get 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.