OpenVX
1. 编译
尝试编译openvx_sample,下载相关代码。
下载的sample code直接使用make可以生成libopenvx.so。
使用python Build.py --os linux可以编译sample code。
2. OpenVX使用流程
主要包含7个部分:
- 创建openvx上下文
vx_context context = vxCreateContext(); - 创建输入、输出图像结点
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 ); - 创建graph
vx_graph graph = vxCreateGraph(context); - 构建graph
vxScaleImageNode(graph, input_rgb_image, output_rgb_image, VX_INTERPOLATION_AREA) - 验证graph
vxVerifyGraph( graph ); - 真正运行graph
vxProcessGraph(graph); - 释放资源
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 {
/*! rief The vx_kernel_e enum */
vx_enum enumeration;
/*! rief The name that kernel will be used with
ef vxGetKernelByName. */
vx_char name[VX_MAX_KERNEL_NAME];
/*! rief The pointer to the function to execute the kernel */
vx_kernel_f function;
/*! rief The pointer to the array of parameter descriptors */
vx_param_description_t *parameters;
/*! rief The number of paraemeters in the array. */
vx_uint32 numParams;
/*! rief The parameters validator */
vx_kernel_validate_f validate;
/*! rief The input validator (deprecated in openvx 1.1) */
void* input_validate;
/*! rief The output validator (deprecated in openvx 1.1) */
void* output_validate;
/*! rief The initialization function */
vx_kernel_initialize_f initialize;
/*! rief 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找不到,直接注释该行代码,程序可以正常运行。但是得到的结果还是不对,全是黑色,好像是没有将处理后的结果拷贝回来,导致结果全是0。
这是因为cl中提供两种形式的表达,一个是image2d_t,一个是简单的buffer,在vx_interface.c中编译cl时,加上了CL_USE_LUMINANCE,使用的是image2d_t;而在编译整个OpenVX时,没有加上CL_USE_LUMINANCE,导致外面使用的是简单的buffer,而一个image2d_t的参数如果使用buffer需要传递5个参数,所以导致最后设置参数时两边不一致出错。修改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));
}
得到的效果正确了,如下:
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