OpenVX

OpenVX

openvx

1. 编译

尝试编译openvx_sample，下载相关代码。

下载的sample code直接使用make可以生成libopenvx.so。

使用python Build.py --os linux可以编译sample code。

2. OpenVX使用流程

主要包含7个部分：

  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。

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

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的相关函数句柄。

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

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

opencl kernel结构:

包含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

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

原文地址：https://www.cnblogs.com/gr-nick/p/9379348.html

时间： 2024-10-08 20:25:29

OpenVX的相关文章

Visionworks OpenVX

[TOC] OpenVX heterogeneous computation framework Spec OpenVX 1.2源碼解析 - 目錄結構 除了官方的參考實作外,下方是不同廠商的實作,有些有開放原始碼有些則是包裝程動態函式庫. Intel Computer Vision SDK AMD OVX : https://github.com/GPUOpen-ProfessionalCompute-Libraries/amdovx-core --> TI OVX: Nvidia Vision

致ubuntu140405：基于window7的vslam开发环境搭建

对于直接安装或解压的程序,除git, cmake, cuda按默认路径安装外,其它程序都安装在D:\app\softname 对于编译安装的程序,都安装在D:\app\softname\build\install 对于ubuntu可首先直接从源安装以下包,之后就只剩下cuda, qt, openni2需要直接安装,其余都编译安装 0.基本开发工具 0.1git&cmake&cuda (1)win安装:安装路径默认,自动追加path环境变量 (2)ubuntu安装:git, cmake3,

基于GPU的图像处理平台

基于GPU的图像处理平台 1.  (309)英伟达推Jetson TX1 GPU模块力推人工智能 1.1 产品概述 Jetson TX1 GPU模块,主要针对近年来蓬勃发展的人工智能市场,包括无人机.机器人等设备. 1.2 处理板技术指标 1． Jetson TX1 GPU模块包括一颗浮点运算达到teraflop级的 2.  基于Maxwell架构的256核心GPU,64位ARM A57芯片组 3.  4GB LPDDR4 RAM内存(每秒带宽速度达25.6GB) 4.  5GB本地存储模块.8

基于GPU的图像处理平台解决方案

基于GPU的图像处理平台 1.  (309)英伟达推Jetson TX1 GPU模块力推人工智能 1.1 产品概述 Jetson TX1 GPU模块,主要针对近年来蓬勃发展的人工智能市场,包括无人机.机器人等设备. 1.2 处理板技术指标 1． Jetson TX1 GPU模块包括一颗浮点运算达到teraflop级的 2.  基于Maxwell架构的256核心GPU,64位ARM A57芯片组 3.  4GB LPDDR4 RAM内存(每秒带宽速度达25.6GB) 4.  5GB本地存储模块.8

NVIDIA Jetson? TX1

NVIDIA® Jetson TX1 是一台模块式计算机,代表了视觉计算领域近20年的研发成就,其尺寸仅有信用卡大小.Jetson TX1 基于崭新 NVIDIA Maxwell™ 架构,配有256个 NVIDIA CUDA® 核心和64位 CPUs,并具有出类拔萃的电源效率.此外,其还囊括了深度学习.计算机视觉.GPU 计算和图像处理等新技术.它是您值得拥有的完善的开发平台——包括开发者组件.工具.开发者平台.社区.支持.培训和产品设计包——因此其可以应付研究.原型设计.开发和生产.总之,Je

NVIDIA Jetson? TX1 Module

NVIDIA® Jetson TX1 是一台模块式计算机,代表了视觉计算领域近20年的研发成就,其尺寸仅有信用卡大小.Jetson TX1 基于NVIDIA Maxwell™ 架构,配有256个 NVIDIA CUDA® 核心和64位 CPUs,并具有出类拔萃的电源效率.此外,其还囊括了深度学习.计算机视觉.GPU 计算和图像处理等新技术.它是您值得拥有的完善的开发平台——包括开发者组件.工具.开发者平台.社区.支持.培训和产品设计包——因此其可以应付研究.原型设计.开发和生产.总之,Jetso

让计算更贴近生活

在这些年,我们伴随着科幻电影作品成长,科幻电影作品中的技术情节让我们激动不已,越来越成熟的高科技产品在不断地改进和创造. 随着计算机技术以及相关辅助学科的发展,人类在高科技领域不断地获得各种飞跃式的突破,科幻电影作品中的场景也不断地被带到现实生 活中来.你是否还记得机器人在进行人物识别过程中的那些场景?机器人对眼前的场景进行扫描,从而来判断眼前的一切是否具有危险存在, 从而作出相关的反映.你是否期望在一个家庭中,有这么一个机器,在你工作下班回家后,其将家庭打扫得一尘不染?...所有的这一切都是

dra7xx: configure for rtos usecase to bulild

example for in hlos add rtos usecase to build. commit f01f061b6db87566e9623ebf43b2466586e1f9fb Date: Fri May 10 17:06:40 2019 +0800 VSDK build: add RTOS usecases to build diff --git a/apps/Makefile b/apps/Makefile index 033b932..03ee1f5 100755 --- a/

原文版：Intel OpenVINO? Toolkit and AWS* Greengrass！！！

https://software.intel.com/en-us/articles/get-started-with-the-openvino-toolkit-and-aws-greengrass Hardware-Accelerated Function-as-a-Service Using AWS Greengrass Hardware Accelerated Function-as-a-Service (FaaS) enables cloud developers to deploy in