OpenVX

时间  2019-11-16
标签 openvx 繁體版
原文   原文链接

OpenVX

openvx 

1. 编译

尝试编译openvx_sample，下载相关代码。
下载的sample code直接使用make能够生成libopenvx.so。
使用python Build.py --os linux能够编译sample code。html

2. OpenVX使用流程

主要包含7个部分：node

  1. 建立openvx上下文
    vx_context context = vxCreateContext();
  2. 建立输入、输出图像结点
    vx_image input_rgb_image = vxCreateImage( context, width, height, VX_DF_IMAGE_RGB );
    vx_image output_rgb_image = vxCreateImage( context, width, height, VX_DF_IMAGE_RGB );
  3. 建立graph
    vx_graph graph = vxCreateGraph(context);
  4. 构建graph
    vxScaleImageNode(graph, input_rgb_image, output_rgb_image, VX_INTERPOLATION_AREA)
  5. 验证graph
    vxVerifyGraph( graph );
  6. 真正运行graph
    vxProcessGraph(graph);
  7. 释放资源
    vxReleaseContext(&context);

3. OpenVX中调用OpenCL代码解析

1. vxCreateContext

一个平台对就一个target，一个target包含多个kernel。python

./sample/framework/vx_context.c中的变量定义了几种target支持, c_model, opencl, openmp：linux

vx_char targetModules[][VX_MAX_TARGET_NAME] = {
    "openvx-c_model",
#if defined(EXPERIMENTAL_USE_OPENCL)
    "openvx-opencl",
#endif
#if defined(EXPERIMENTAL_USE_OPENMP)
    "openvx-openmp"
#endif
};

以OpenCL为例，当用户调用函数vxCreateContext(sample/framework/vx_context.c)时，其会调用函数ownLoadTarget (sample/framework/vx_target.c), 去dlopen打开libopenvx-opencl.so, 使用dlsym(mod, name)获取vxTargetInit, vxTargetAddKernel(sample/targets/opencl/vx_interface.c)等opencl的相关函数句柄。git

而在vxTargetAddKernel函数中，调用ownInitializeKernel(sample/framework/vx_kernel.c)加载了全部OpenCL实现的kernel函数。github

在sample/targets/opencl目录下的c文件定义了一些vx_cl_kernel_description_t box3x3_clkernel变量，包括box3x3_clkernel, gaussian3x3_clkernel, and_kernel等 ，这些kernelapi

opencl kernel结构:ide

包含vx_kernel_description_t还有一些其它属性，它把function置为NULL，并提供了一个sourcepath变量用来存放opencl函数。函数

typedef struct _vx_cl_kernel_description_t {
    vx_kernel_description_t description;
    char             sourcepath[VX_CL_MAX_PATH];
    char             kernelname[VX_MAX_KERNEL_NAME];
    cl_program       program[VX_CL_MAX_PLATFORMS];
    cl_kernel        kernels[VX_CL_MAX_PLATFORMS];
    cl_uint          num_kernels[VX_CL_MAX_PLATFORMS];
    cl_int           returns[VX_CL_MAX_PLATFORMS][VX_CL_MAX_DEVICES];
    void            *reserved; /* for additional data */
} vx_cl_kernel_description_t;

kernel结构:学习

typedef struct _vx_kernel_description_t {
    /*! \brief The vx_kernel_e enum */
    vx_enum                 enumeration;
    /*! \brief The name that kernel will be used with \ref vxGetKernelByName. */
    vx_char                 name[VX_MAX_KERNEL_NAME];
    /*! \brief The pointer to the function to execute the kernel */
    vx_kernel_f             function;
    /*! \brief The pointer to the array of parameter descriptors */
    vx_param_description_t *parameters;
    /*! \brief The number of paraemeters in the array. */
    vx_uint32               numParams;
    /*! \brief The parameters validator */
    vx_kernel_validate_f    validate;
    /*! \brief The input validator (deprecated in openvx 1.1) */
    void* input_validate;
    /*! \brief The output validator (deprecated in openvx 1.1) */
    void* output_validate;
    /*! \brief The initialization function */
    vx_kernel_initialize_f initialize;
    /*! \brief The deinitialization function */
    vx_kernel_deinitialize_f deinitialize;
} vx_kernel_description_t;

能够看到目前虽然配置了一些参数，但OpenCL分为主机端代码和device端代码，device端代码在kernel/opencl中，而host端代码在哪呢？如何根据设置的参数去执行Host端代码，从而执行device端代码：
能够看到在vxTargetInit函数中，调用ownInitializeKernel初始化kernel时，判断了kfunc是否为NULL，(kfunc == NULL ? vxclCallOpenCLKernel : kfunc)若是为NULL则使用vxclCallOpenCLKernel函数。

咱们再看vxclCallOpenCLKernel函数，咱们发现这个函数里有clSetKernelArg，clEnqueueNDRangeKernel等OpenCL的API函数，这个即是host-side的OpenCL代码。

2. vxScaleImageNode

在sample/framework/vx_node_api.c中定义了全部提供的可用的OpenVX结点，包括vxScaleImageNode结点，经过以下方法建立Node：

vx_kernel kernel   = vxGetKernelByEnum( context, VX_KERNEL_SCALE_IMAGE );

若是函数有两种实现，那么按照优先级使用： opencl > openmp > c_model。（不对，感受优先使用的是c_model的函数；实际是先找到opencl kernel，但找到以后并无中止查找，找到后面的c_model就会覆盖掉前面的opencl kernel。不知道这儿是写错了，仍是就是要优先使用c_model，代码见sample/framework/vx_kernel.c中的vxGetKernelByEnum函数）

node的参数如何传递给kernel： 在vxCreateNodeByStructure中调用vxSetParameterByIndex将Node的参数传递kernel。

3. vxVerifyGraph

vx_graph.c会调用每个结点的validator函数，包括inputValidator，outputValidator，确保构建的Graph能够跑通。

4. vxProcessGraph

vxProcessGraph函数调用vxExecuteGraph函数，在其中调用action = target->funcs.process(target, &node, 0, 1);，其中的funcs.process就是各个target的vxTargetProcess函数。

在vxTargetProcess中会调用nodes[n]->kernel->function，即咱们事先定义的host-side端代码，传递结点，参数，以及参数个数：

status = nodes[n]->kernel->function((vx_node)nodes[n], 
                                            (vx_reference *)nodes[n]->parameters,
                                             nodes[n]->kernel->signature.num_parameters);

而咱们的function，则主要负责内存管理，以及调用device端代码。

几种参数类型：
memory:
CL_MEM_OBJECT_BUFFER
CL_MEM_OBJECT_IMAGE2D
scalar:
VX_TYPE_SCALAR
threashold:
VX_TYPE_THRESHOLD

4. OpenVX中使用OpenCL的编译问题

使用Makefile编译出来的so默认是没有opencl。

使用Build.py出来的so能够有opencl，但结点报错：
Target[1] is not valid!
Target[2] is not valid!
LOG: [ status = -17 ] Node: org.khronos.openvx.color_convert: parameter[1] is not a valid type 1280!

在target.mak中对SYSDEFS添加EXPERIMENTAL_USE_OPENCL，能够编译Opencl，但在运行时build opencl 代码时报错，能够将错误信息打印出来，发现找不到头文件。

查看代码，发如今sample/targets/opencl/vx_interface.c中须要以下两个参数，VX_CL_INCLUDE_DIR是VX头文件位置，VX_CL_SOURCE_DIR是CL源码位置，在环境中能够配置这两个参数：
char *vx_incs = getenv("VX_CL_INCLUDE_DIR");
char *cl_dirs = getenv("VX_CL_SOURCE_DIR");

/usr/include/features.h:367:12: fatal error: 'sys/cdefs.h' file not found
在cl编译命令里(sample/targets/opencl/vx_interface.c)添加-I /usr/include/x86_64-linux-gnu/:
snprintf(cl_args, sizeof(cl_args), "-D VX_CL_KERNEL -I %s -I /usr/include/x86_64-linux-gnu/ -I %s %s %s", vx_incs, cl_dirs...

Linux gnu/stubs-32.h: No such file or directory
这是缺乏32位的嵌入式C库。在嵌入式开发环境配置时，也常遇到这个问题。sudo apt-get install libc6-dev-i386

fatal error: 'stddef.h' file not found
定位stddef.h, 在cl编译命令里cl_args里添加-I /usr/include/linux/。

vx_khr_opencl.h和vx_api.h里有些类型进行了重定义：
不要在vx_khr_opencl.h里include vx_api.h。

histogram.cl仍然报错，将histogram的kernel去掉，就能够成功编译。

5. 使用OpenCL vx_not

使用c_model的VX_KERNEL_NOT能够正常运行，使用opencl的就会报以下错误：

clSetKernelArg: OpenCL error CL_INVALID_ARG_INDEX at vxclCallOpenCLKernel in /home/guru_ge/openvx/openvx_sample/sample/targets/opencl/vx_interface.c:639
clSetKernelArg: OpenCL error CL_INVALID_ARG_INDEX at vxclCallOpenCLKernel in /home/guru_ge/openvx/openvx_sample/sample/targets/opencl/vx_interface.c:639
clEnqueueNDRangeKernel: OpenCL error CL_INVALID_KERNEL_ARGS at vxclCallOpenCLKernel in /home/guru_ge/openvx/openvx_sample/sample/targets/opencl/vx_interface.c:724

尝试本身写host side的code。

写完发现并在Load时就不经过，检查缘由，打开log信息，发如今本身实现的代码中有个CL_ERROR_MSG找不到，直接注释该行代码，程序能够正常运行。可是获得的结果仍是不对，全是黑色，好像是没有将处理后的结果拷贝回来，致使结果全是０。

这是由于cl中提供两种形式的表达，一个是image2d_t，一个是简单的buffer，在vx_interface.c中编译cl时，加上了CL_USE_LUMINANCE，使用的是image2d_t；而在编译整个OpenVX时，没有加上CL_USE_LUMINANCE，致使外面使用的是简单的buffer，而一个image2d_t的参数若是使用buffer须要传递５个参数，因此致使最后设置参数时两边不一致出错。修改concerto/target.mak在31行SYSDEFS里加上CL_USE_LUMINANCE就能够了。

虽然不报错了，可是出来的结果竟然是一条直线，而不是取反后的效果，很奇怪：

一条直线
一条直线

难道是传给opencl的图像就不对？尝试手动拷贝图像数据。

尝试学习opencl c 语法，修改代码查看结果，发现openvx在实现opencl的时候not kernel时存在一些不规范的地方，可能这些问题在其它平台能够运行，但到如今这个平台上就不行了。

原来的kernel实现：

__kernel void vx_not(read_only image2d_t a, write_only image2d_t b) {
    int2 coord = (get_global_id(0), get_global_id(1));
    write_imageui(b, coord, ~read_imageui(a, nearest_clamp, coord));
}

首先我尝试打印其像素坐标时，发现获得的x, y坐标老是相同的，这很奇怪，这也解释了为何结果只有一条直线，由于它只写了x, y坐标相同的那些像素点的值。查看 API发现get_global_id返回的是size_t，因此要用(int)去显示转换一下，再打印时，发现坐标在不停的变换，变成正常的了。

再运行，获得的图竟然是一幅全白的图，说明像素值还有问题。尝试打印原像素值，与取反后的像素值，发现相加不是255，说明这里的取反操做也有问题。read_imageui返回的类型是uint4向量，咱们取反时，获得的结果并不对，这里使用255直接相减，最后代码以下所示：

__kernel void vx_not(read_only image2d_t a, write_only image2d_t b) {
    int2 coord = (int2)(get_global_id(0), get_global_id(1));
    write_imageui(b, coord, 255-read_imageui(a, nearest_clamp, coord));
}

获得的效果正确了，以下：

right_result
right_result

6. 实现OpenCL vx_scale

实现opencl scale报错：
parameter[1] is an invalid dimension 640x240

传递的参数是(inputImg, outputImg, type)，parameter[1]应该是输出图像，大小确实应该是640x240。
使用c_model中的outputvalidator就不报这个错了，说明不能直接return VX_SUCCESS，可能validator中还须要作些其它的事情。
在validator中会记录一些信息，以供后面verify时与实际传入参数比对，因此不能直接返回SUCCESS:

ptr->type = VX_TYPE_IMAGE;
    ptr->dim.image.format = VX_DF_IMAGE_U8;                         
    ptr->dim.image.width = width;                                   
    ptr->dim.image.height = height;

然而如今又报以下错误：
clEnqueueNDRangeKernel: OpenCL error CL_INVALID_EVENT at vxclCallOpenCLKernel in /home/guru_ge/openvx/openvx_sample/sample/targets/opencl/vx_interface.c:725
clEnqueueReadImage: OpenCL error CL_INVALID_EVENT at vxclCallOpenCLKernel in /home/guru_ge/openvx/openvx_sample/sample/targets/opencl/vx_interface.c:793

为何event会invalid呢？尝试本身写host-side代码，不使用默认的。

本身的代码报以下错误：
VX_ZONE_ERROR:[vxcl_platform_notifier:59] CL_OUT_OF_RESOURCES error executing CL_COMMAND_READ_IMAGE on GeForce GTX 1080 Ti (Device 0)

spec里解释CL_OUT_OF_RESOURCES: if there is a failure to allocate resources required by the OpenCL implementation on the device.

这估计是使用c_model的validator致使没有初始化cl_mem，尝试使用cl validator。

在check scale node parameter时报以下错误：

LOG: [ status = -10 ] Node[3] org.khronos.openvx.image_scaling: parameter[2] failed input/bi validation!

这估计是Input validator里只容许Image类型，没有判断scalar类型。

因此validator要对每一个参数逐一判断，对于input参数，直接返回SUCCESS就能够了；而对output参数，还须要写一些信息。

结果仍是全黑的，在kernel中打印坐标发现也不对，查看代码发现输入的维度是输入图片的大小，这儿应该是输出图像的大小才对。

再运行仍是黑色，发如今取坐标转换时，没有将float转为int，致使有问题（因此类型要确保彻底一致，不会替你作转换）。修改后，能够正常运行。

__kernel void image_scaling(read_only image2d_t in, write_only image2d_t out) {
    //从glob_id中获取目标像素坐标
    int2 coordinate = (int2)(get_global_id(0), get_global_id(1));
    //计算归一化浮点坐标 
    float2 normalizedCoordinate = convert_float2(coordinate) * (float2)(2, 2);
    //根据归一化坐标从原图中读取像素数据
    uint4 colour = read_imageui(in, sampler, convert_int2(normalizedCoordinate));
    //将像素数据写入目标图像 
    write_imageui(out, coordinate, colour);
}

实际比较，vx_not, vx_scale使用opencl, c_model实现时间对比：
opencl:
average time: 44099.857143 us

c_model:
average time: 68343.380952 us

7. vx debug print信息

程序中经过获取VX_ZONE_MASK环境变量的值来设置Log级别，能够经过以下将全部级别信息都打开：
export VX_ZONE_MASK=fffff

一共有以下几个级别，每一个级别占int的一个bit位：

enum vx_debug_zone_e {
    VX_ZONE_ERROR       = 0,    /*!< Used for most errors */
    VX_ZONE_WARNING     = 1,    /*!< Used to warning developers of possible issues */
    VX_ZONE_API         = 2,    /*!< Used to trace API calls and return values */
    VX_ZONE_INFO        = 3,    /*!< Used to show run-time processing debug */

    VX_ZONE_PERF        = 4,    /*!< Used to show performance information */
    VX_ZONE_CONTEXT     = 5,
    VX_ZONE_OSAL        = 6,
    VX_ZONE_REFERENCE   = 7,

    VX_ZONE_ARRAY       = 8,
    VX_ZONE_IMAGE       = 9,
    VX_ZONE_SCALAR      = 10,
    VX_ZONE_KERNEL      = 11,

    VX_ZONE_GRAPH       = 12,
    VX_ZONE_NODE        = 13,
    VX_ZONE_PARAMETER   = 14,
    VX_ZONE_DELAY       = 15,

    VX_ZONE_TARGET      = 16,
    VX_ZONE_LOG         = 17,

    VX_ZONE_MAX         = 32
};

Ref

AMD openvx实现：

https://github.com/GPUOpen-ProfessionalCompute-Libraries/amdovx-core
https://github.com/GPUOpen-ProfessionalCompute-Libraries/amdovx-modules

相关文章
相关标签/搜索
每日一句
    每一个你不满意的现在，都有一个你没有努力的曾经。
本站公众号
   欢迎关注本站公众号,获取更多信息
联系我们 最近搜索 最新文章 沪ICP备13005482号-10 MyBatis教程 SQL 教程 MySQL教程 Java 教程 Thymeleaf 教程 Hibernate教程 Spring教程 Redis教程