第一节 openvx(PC)

安装相关软件：

1. 下载VivanteIDE 2.1.0软件

    /buildroot/output/build/gpu/gpu_utils/VivantelIDE

备注：openvx的API可在/buildroot/output/build/gpu-1.0.0/gpu_driver/driver/khronos/libOpenvx/utility/vx_utility.c

2. 安装VivanteIDE 2.1.0软件

 备注：在安装过程中一般选择默认安装就好，切记不要包含空格和中文, 否则会出错

 第一步： 调用OpenVX官方手册中的Node使用示例

1. 打开VivantelIDE2.1.0开发软件, 新建一个自定义文件夹E:\Program\OpenVX\test

2. 将安装路径下的C:\VeriSilicon\VivanteIDE2.1.0\examples(根据自己的安装目录来定)中的文件夹vx_tutorial2拷贝到

E:\OpenVX\test下，得到如下结果:

3. 右击Project Explorer的空白处，在弹出的选项中选择ImportProject选项卡的General选项卡中选择Existing Projects intoWorkspace选项.

4. 将E:\OpenVX\test\vx_tutorial2工程导入到当前的Workspace. 单击finish按钮完成导入.

5. 选中vx_tutorial2，右击选择Build Project编译整个工程.

6. 选中vx_tutorial2，右击选择RunAsàLocal Old C/C++ Application(一定是这个选择项)就可以执行整个工程。

   

7. 出现vx process success，表示运行成功.(输入lena_gray.bmp, 输出lena_sobel.bmp)

第二部分：自定义Node示例

a). 创建自定义的Kernel程序

1.  点击FileàNew选择OpenVX C++ Project,如果没有弹出需要自己在C/C++选项中寻找。

2. 接着在弹出的项目设置窗口中设置项目的名字(ImageAdder),项目的类型必须选择Shared Library,接着自定义配置一下项目作者及所有权选项.

3. 在跳出的定义设置窗口需要设置kernel的枚举值以及后缀名称

4. 接下来配置这个Kernel的参数个数，类型以及方向。这里我们是两个Image(short类型)输入,一个Image(short类型)输出

5. 设置完成后，我们可以直接点击finish完成.

  此时 E:\OpenVX\test多出ImageAdder文件夹, ImageAdder文件夹下包含以下内容:

  

6.  将ImageAdder.cpp中的vxImageAdderValidator函数中的VX_IMAGE_FORMAT 图形格式更改为VX_DF_IMAGE_S16. 否则参数验证出错.

       if (format != VX_DF_IMAGE_U16)----------  if (format != VX_DF_IMAGE_S16)

       if (format != VX_DF_IMAGE_U16)----------  if (format != VX_DF_IMAGE_S16)

       vx_df_image format = VX_DF_IMAGE_U8;---------vx_df_image format = VX_DF_IMAGE_S16;

7. 点击项目中的ImageAdder，编译ImageAdder, 在Debug下生成动态库ImageAdder.dll, 这个dll供vx_tutorial2.c使用.

  

b). 创建C++应用程序

1. vx_tutorial2.c中书写程序.

/*
 *
 * Created on：2018.3.29
 *    Author: WangZhen
 */
#include <VX/vx.h>
#include <VX/vxu.h>
#include <VX/vx_lib_extras.h>
#include <VX/vx_api.h>
#include <VX/vx_khr_cnn.h>

#include <stdio.h>
#include <string.h>

#include "../ImageAdder/vx_lib_imageadder.h"

#define OBJCHECK(objVX) if(!(objVX)) { printf("[%s : %d] %s\n",__FILE__, __LINE__, "obj create error.");return -1; }
#define FUNCHECK(funRet) if(VX_SUCCESS!=(funRet)) { printf("[%s : %d] %s\n",__FILE__, __LINE__, "function error.");vxReleaseContext(&context);return -1;}

#define WIDTH  33
#define HEIGHT 16

#define SIZE (WIDTH*HEIGHT)

int main()
{
	/* VX variables */
	vx_status status = VX_FAILURE;
	vx_context context = NULL;
	vx_graph graph = NULL;
	vx_node node = NULL;

	 /* vx procedure */
	context = vxCreateContext();
    OBJCHECK(context);
    graph = vxCreateGraph(context);
    OBJCHECK(graph);

    /* read data */
    short *buf_a = (short*)malloc(SIZE*sizeof(short));
    short *buf_b = (short*)malloc(SIZE*sizeof(short));
    short *buf_c = (short*)malloc(SIZE*sizeof(short));
    for(int i=0;i<SIZE;i++)
    {
    	buf_a[i] = i;
    	buf_b[i] = i+1;
    	buf_c[i] = buf_a[i] + buf_b[i];
    }

    vx_image A = vxCreateImage(context,WIDTH,HEIGHT,VX_DF_IMAGE_S16);
    OBJCHECK(A);
    vx_image B = vxCreateImage(context,WIDTH,HEIGHT,VX_DF_IMAGE_S16);
    OBJCHECK(B);
    vx_image C = vxCreateImage(context,WIDTH,HEIGHT,VX_DF_IMAGE_S16);
    OBJCHECK(C);

    /* transfer image from cpu to gpu */
    vx_imagepatch_addressing_t imgInfo = VX_IMAGEPATCH_ADDR_INIT;
    short* imgData = 0;
    vx_rectangle_t rect = {0,0,WIDTH,HEIGHT};
    vx_map_id map_id = 0;

    status = vxMapImagePatch(A,&rect,0,&map_id,&imgInfo,(void**)&imgData,VX_WRITE_ONLY,VX_MEMORY_TYPE_HOST,0);
    FUNCHECK(status);
    memcpy(imgData,buf_a,sizeof(vx_int16)*SIZE);
    status = vxUnmapImagePatch(A,map_id); //match with vxUnmapImagePatch()
    FUNCHECK(status);
    imgData = NULL;

    status = vxMapImagePatch(B,&rect,0,&map_id,&imgInfo,(void**)&imgData,VX_WRITE_ONLY,VX_MEMORY_TYPE_HOST,0);
    FUNCHECK(status);
    memcpy(imgData,buf_b,sizeof(vx_int16)*SIZE);
    status = vxUnmapImagePatch(B,map_id);
    FUNCHECK(status);
    imgData = NULL;

    /* process image */
    printf("before");

    node = vxcImageAdderNode(graph,A,B,C);
    OBJCHECK(node);

    status = vxVerifyGraph(graph);
    FUNCHECK(status);

    status = vxProcessGraph(graph);
    FUNCHECK(status);

    /* transfer image from gpu to cpu */
    status = vxMapImagePatch(C,&rect,0,&map_id,&imgInfo,(void**)&imgData,VX_READ_ONLY,VX_MEMORY_TYPE_HOST,0);
    FUNCHECK(status);

    int errNum = 0;
    for(int i=0;i<SIZE;i++)
    {
    	if(buf_c[i] != imgData[i])
    	{
    		printf("not match:[%d] %d vs %\n",i,buf_c[i],imgData[i]);
    		errNum++;
    	}
    }
    if(errNum == 0)
    {
    	printf("test pass\n");
    }

    status = vxUnmapImagePatch(C,map_id);
    FUNCHECK(status);
    imgData = NULL;

    free(buf_a);
    free(buf_b);
    free(buf_c);
    vxReleaseContext(&context);

	return 0;
}

2. 将kernel文件生成的ImageAdder.dll文件拷贝到当前工程目录中，注意

    vx_tutorial2.c中的 node = vxcImageAdderNode(graph,A,B,C); 其中的名称必须与vx_lib_imageadder.h中

    的node名称相一致.

3. 这时运行vx_tutorial2这个工程发现运算结果错误，原因是我们并没有编写ImageAdder.vx这个文件的代码.

4. ImageAdder.cpp中的参数先不用修改，先跑通程序先.

5. 书写ImageAdder_vx.vx代码

/*
 ============================================================================
 Name        : ImageAdder.vx
 Author      : wangzhen
 Version     :
 Copyright   : 
 Description : 
 ============================================================================
 */
#include "cl_viv_vx_ext.h"
_viv_uniform VXC_512Bits UniDP2x8_ADD;

__kernel void ImageAdderVXC
    (
	image2d_t InputA,
	image2d_t InputB,
	image2d_t Output
    )
{
	int2 coord = (int2)(get_global_id(0),get_global_id(1)); //其中get_global_id(0)/(1)跟global work scale息息相关，图像分块的概念，每个线程从哪里开始运行kernel.
	vxc_short8 inputA,inputB,output; //vxc_short8, 类似数组概念,  8个short数据类型. 一个线程一次处理8个short数据.
	VXC_ReadImage(inputA, InputA, coord, VXC_5BITOFFSET_XY(0,0), VXC_MODIFIER(0,7, 0, VXC_RM_TowardZero, 0)); //从InputA中一次读取8个数据inputA
	VXC_ReadImage(inputB, InputB, coord, VXC_5BITOFFSET_XY(0,0), VXC_MODIFIER(0,7, 0, VXC_RM_TowardZero, 0)); //从InputB中一次读取8个数据inputB
	VXC_DP2x8(output, inputA, inputB, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 1), UniDP2x8_ADD01); //DP指令处理inputA+inputB的操作
	VXC_WriteImage(Output, coord, output, VXC_MODIFIER(0, 7, 0, VXC_RM_TowardZero, 0)); //将output数据写回到Output中
}

6. 打开DP指令软件，对DP指令Unidp_16ShortAdder在EVIS编辑器的代码如下：

    New: 新建

    Open: 打开

    Save: 保存

    EVIS1: VIP7000, EVIS2：VIP8000

    DP2x8: 一个kernel，一个线程. 比如一个线程处理x方向8个像素，y方向1个像素. 在一个线程内部实现a + b这个操作，

                 一个a, 一个b,有2个(控制config的个数), 故为2.

                 一个线程处理x方向8个像素，y方向1个像素，故输出为8个结果, 故为8

    AccumType: 多少个bin看做一个数据整体.

                         比如选择S16 short16位, 也即2个bin看做一个数据. 选择S32 short32位, 也即4个bin看做一个整体.

   Gene: 生成DP代码.

   Constant: 系数, eg: S8, 系数范围[-128,127].

  SRC0: 寄存器, DP指令代码中参数inputA的数据类型

  SRC1: 寄存器, DP指令代码中参数inputB的数据类型

7. 点击Gene按钮，将生成的类似汇编的代码如下图所示：

  

  8. 将生成的代码拷贝至ImageAdder.cpp中的vxImageAdderInitializer()，代码如下：

      9. 根据图像尺寸等具体情况，修改shaderParam的参数.

       

10.  再次运行vx_tutorail2工程，可以看到控制台输出tess pass的字样，表示程序运行成功.

      备注：里面的参数不一定是最佳，但是可以确保程序运行成功.

                openvx的API函数，参见OpenVX_Specification_1_1.pdf.