OpenVX
OpenVX
openvx
1. 編譯
嘗試編譯openvx_sample，下載相關(guān)代碼。
下載的sample code直接使用make可以生成libopenvx.so。
使用python Build.py --os linux可以編譯sample code。
2. OpenVX使用流程
主要包含7個(gè)部分：
創(chuàng)建openvx上下文
vx_context context = vxCreateContext();
創(chuàng)建輸入、輸出圖像結(jié)點(diǎn)
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 );
創(chuàng)建graph
vx_graph graph = vxCreateGraph(context);
構(gòu)建graph
vxScaleImageNode(graph, input_rgb_image, output_rgb_image, VX_INTERPOLATION_AREA)
驗(yàn)證graph
vxVerifyGraph( graph );
真正運(yùn)行g(shù)raph
vxProcessGraph(graph);
釋放資源
vxReleaseContext(&context);
3. OpenVX中調(diào)用OpenCL代碼解析
1. vxCreateContext
一個(gè)平臺(tái)對(duì)就一個(gè)target，一個(gè)target包含多個(gè)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
};
以O(shè)penCL為例，當(dāng)用戶調(diào)用函數(shù)vxCreateContext(sample/framework/vx_context.c)時(shí)，其會(huì)調(diào)用函數(shù)ownLoadTarget (sample/framework/vx_target.c), 去dlopen打開(kāi)libopenvx-opencl.so, 使用dlsym(mod, name)獲取vxTargetInit, vxTargetAddKernel(sample/targets/opencl/vx_interface.c)等opencl的相關(guān)函數(shù)句柄。
而在vxTargetAddKernel函數(shù)中，調(diào)用ownInitializeKernel(sample/framework/vx_kernel.c)加載了所有OpenCL實(shí)現(xiàn)的kernel函數(shù)。
在sample/targets/opencl目錄下的c文件定義了一些vx_cl_kernel_description_t box3x3_clkernel變量，包括box3x3_clkernel, gaussian3x3_clkernel, and_kernel等 ，這些kernel
opencl kernel結(jié)構(gòu):
包含vx_kernel_description_t還有一些其它屬性，它把function置為NULL，并提供了一個(gè)sourcepath變量用來(lái)存放opencl函數(shù)。
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結(jié)構(gòu):
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;
可以看到目前雖然配置了一些參數(shù)，但OpenCL分為主機(jī)端代碼和device端代碼，device端代碼在kernel/opencl中，而host端代碼在哪呢？如何根據(jù)設(shè)置的參數(shù)去執(zhí)行Host端代碼，從而執(zhí)行device端代碼：
可以看到在vxTargetInit函數(shù)中，調(diào)用ownInitializeKernel初始化kernel時(shí)，判斷了kfunc是否為NULL，(kfunc == NULL ? vxclCallOpenCLKernel : kfunc)如果為NULL則使用vxclCallOpenCLKernel函數(shù)。
我們?cè)倏磛xclCallOpenCLKernel函數(shù)，我們發(fā)現(xiàn)這個(gè)函數(shù)里有clSetKernelArg，clEnqueueNDRangeKernel等OpenCL的API函數(shù)，這個(gè)便是host-side的OpenCL代碼。
2. vxScaleImageNode
在sample/framework/vx_node_api.c中定義了所有提供的可用的OpenVX結(jié)點(diǎn)，包括vxScaleImageNode結(jié)點(diǎn)，通過(guò)如下方法創(chuàng)建Node：
vx_kernel kernel = vxGetKernelByEnum( context, VX_KERNEL_SCALE_IMAGE );
如果函數(shù)有兩種實(shí)現(xiàn)，那么按照優(yōu)先級(jí)使用： opencl > openmp > c_model。（不對(duì)，感覺(jué)優(yōu)先使用的是c_model的函數(shù)；實(shí)際是先找到opencl kernel，但找到之后并沒(méi)有停止查找，找到后面的c_model就會(huì)覆蓋掉前面的opencl kernel。不知道這兒是寫(xiě)錯(cuò)了，還是就是要優(yōu)先使用c_model，代碼見(jiàn)sample/framework/vx_kernel.c中的vxGetKernelByEnum函數(shù)）
node的參數(shù)如何傳遞給kernel： 在vxCreateNodeByStructure中調(diào)用vxSetParameterByIndex將Node的參數(shù)傳遞kernel。
3. vxVerifyGraph
vx_graph.c會(huì)調(diào)用每一個(gè)結(jié)點(diǎn)的validator函數(shù)，包括inputValidator，outputValidator，確保構(gòu)建的Graph可以跑通。
4. vxProcessGraph
vxProcessGraph函數(shù)調(diào)用vxExecuteGraph函數(shù)，在其中調(diào)用action = target->funcs.process(target, &node, 0, 1);，其中的funcs.process就是各個(gè)target的vxTargetProcess函數(shù)。
在vxTargetProcess中會(huì)調(diào)用nodes[n]->kernel->function，即我們事先定義的host-side端代碼，傳遞結(jié)點(diǎn)，參數(shù)，以及參數(shù)個(gè)數(shù)：
status = nodes[n]->kernel->function((vx_node)nodes[n],
(vx_reference *)nodes[n]->parameters,
nodes[n]->kernel->signature.num_parameters);
而我們的function，則主要負(fù)責(zé)內(nèi)存管理，以及調(diào)用device端代碼。
幾種參數(shù)類型：
memory:
CL_MEM_OBJECT_BUFFER
CL_MEM_OBJECT_IMAGE2D
scalar:
VX_TYPE_SCALAR
threashold:
VX_TYPE_THRESHOLD
4. OpenVX中使用OpenCL的編譯問(wèn)題
使用Makefile編譯出來(lái)的so默認(rèn)是沒(méi)有opencl。
使用Build.py出來(lái)的so可以有opencl，但結(jié)點(diǎn)報(bào)錯(cuò)：
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中對(duì)SYSDEFS添加EXPERIMENTAL_USE_OPENCL，可以編譯Opencl，但在運(yùn)行時(shí)build opencl 代碼時(shí)報(bào)錯(cuò)，可以將錯(cuò)誤信息打印出來(lái)，發(fā)現(xiàn)找不到頭文件。
查看代碼，發(fā)現(xiàn)在sample/targets/opencl/vx_interface.c中需要如下兩個(gè)參數(shù)，VX_CL_INCLUDE_DIR是VX頭文件位置，VX_CL_SOURCE_DIR是CL源碼位置，在環(huán)境中可以配置這兩個(gè)參數(shù)：
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庫(kù)。在嵌入式開(kāi)發(fā)環(huán)境配置時(shí)，也常遇到這個(gè)問(wèn)題。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里有些類型進(jìn)行了重定義：
不要在vx_khr_opencl.h里include vx_api.h。
histogram.cl仍然報(bào)錯(cuò)，將histogram的kernel去掉，就可以成功編譯。
5. 使用OpenCL vx_not
使用c_model的VX_KERNEL_NOT可以正常運(yùn)行，使用opencl的就會(huì)報(bào)如下錯(cuò)誤：
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
嘗試自己寫(xiě)host side的code。
寫(xiě)完發(fā)現(xiàn)并在Load時(shí)就不通過(guò)，檢查原因，打開(kāi)log信息，發(fā)現(xiàn)在自己實(shí)現(xiàn)的代碼中有個(gè)CL_ERROR_MSG找不到，直接注釋該行代碼，程序可以正常運(yùn)行。但是得到的結(jié)果還是不對(duì)，全是黑色，好像是沒(méi)有將處理后的結(jié)果拷貝回來(lái)，導(dǎo)致結(jié)果全是０。
這是因?yàn)閏l中提供兩種形式的表達(dá)，一個(gè)是image2d_t，一個(gè)是簡(jiǎn)單的buffer，在vx_interface.c中編譯cl時(shí)，加上了CL_USE_LUMINANCE，使用的是image2d_t；而在編譯整個(gè)OpenVX時(shí)，沒(méi)有加上CL_USE_LUMINANCE，導(dǎo)致外面使用的是簡(jiǎn)單的buffer，而一個(gè)image2d_t的參數(shù)如果使用buffer需要傳遞５個(gè)參數(shù)，所以導(dǎo)致最后設(shè)置參數(shù)時(shí)兩邊不一致出錯(cuò)。修改concerto/target.mak在31行SYSDEFS里加上CL_USE_LUMINANCE就可以了。
雖然不報(bào)錯(cuò)了，但是出來(lái)的結(jié)果居然是一條直線，而不是取反后的效果，很奇怪：
一條直線
難道是傳給opencl的圖像就不對(duì)？嘗試手動(dòng)拷貝圖像數(shù)據(jù)。
嘗試學(xué)習(xí)opencl c 語(yǔ)法，修改代碼查看結(jié)果，發(fā)現(xiàn)openvx在實(shí)現(xiàn)opencl的時(shí)候not kernel時(shí)存在一些不規(guī)范的地方，可能這些問(wèn)題在其它平臺(tái)可以運(yùn)行，但到現(xiàn)在這個(gè)平臺(tái)上就不行了。
原來(lái)的kernel實(shí)現(xiàn)：
__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));
}
首先我嘗試打印其像素坐標(biāo)時(shí)，發(fā)現(xiàn)得到的x, y坐標(biāo)總是相同的，這很奇怪，這也解釋了為什么結(jié)果只有一條直線，因?yàn)樗粚?xiě)了x, y坐標(biāo)相同的那些像素點(diǎn)的值。查看 API發(fā)現(xiàn)get_global_id返回的是size_t，所以要用(int)去顯示轉(zhuǎn)換一下，再打印時(shí)，發(fā)現(xiàn)坐標(biāo)在不停的變換，變成正常的了。
再運(yùn)行，得到的圖居然是一幅全白的圖，說(shuō)明像素值還有問(wèn)題。嘗試打印原像素值，與取反后的像素值，發(fā)現(xiàn)相加不是255，說(shuō)明這里的取反操作也有問(wèn)題。read_imageui返回的類型是uint4向量，我們?nèi)》磿r(shí)，得到的結(jié)果并不對(duì)，這里使用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. 實(shí)現(xiàn)OpenCL vx_scale
實(shí)現(xiàn)opencl scale報(bào)錯(cuò)：
parameter[1] is an invalid dimension 640x240
傳遞的參數(shù)是(inputImg, outputImg, type)，parameter[1]應(yīng)該是輸出圖像，大小確實(shí)應(yīng)該是640x240。
使用c_model中的outputvalidator就不報(bào)這個(gè)錯(cuò)了，說(shuō)明不能直接return VX_SUCCESS，可能validator中還需要做些其它的事情。
在validator中會(huì)記錄一些信息，以供后面verify時(shí)與實(shí)際傳入?yún)?shù)比對(duì)，所以不能直接返回SUCCESS:
ptr->type = VX_TYPE_IMAGE;
ptr->dim.image.format = VX_DF_IMAGE_U8;
ptr->dim.image.width = width;
ptr->dim.image.height = height;
然而現(xiàn)在又報(bào)如下錯(cuò)誤：
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會(huì)invalid呢？嘗試自己寫(xiě)host-side代碼，不使用默認(rèn)的。
自己的代碼報(bào)如下錯(cuò)誤：
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.
這估計(jì)是使用c_model的validator導(dǎo)致沒(méi)有初始化cl_mem，嘗試使用cl validator。
在check scale node parameter時(shí)報(bào)如下錯(cuò)誤：
LOG: [ status = -10 ] Node[3] org.khronos.openvx.image_scaling: parameter[2] failed input/bi validation!
這估計(jì)是Input validator里只允許Image類型，沒(méi)有判斷scalar類型。
所以validator要對(duì)每個(gè)參數(shù)逐一判斷，對(duì)于input參數(shù)，直接返回SUCCESS就可以了；而對(duì)output參數(shù)，還需要寫(xiě)一些信息。
結(jié)果還是全黑的，在kernel中打印坐標(biāo)發(fā)現(xiàn)也不對(duì)，查看代碼發(fā)現(xiàn)輸入的維度是輸入圖片的大小，這兒應(yīng)該是輸出圖像的大小才對(duì)。
再運(yùn)行還是黑色，發(fā)現(xiàn)在取坐標(biāo)轉(zhuǎn)換時(shí)，沒(méi)有將float轉(zhuǎn)為int，導(dǎo)致有問(wèn)題（所以類型要確保完全一致，不會(huì)替你做轉(zhuǎn)換）。修改后，可以正常運(yùn)行。
__kernel void image_scaling(read_only image2d_t in,
write_only image2d_t out)
{
//從glob_id中獲取目標(biāo)像素坐標(biāo)
int2 coordinate = (int2)(get_global_id(0), get_global_id(1));
//計(jì)算歸一化浮點(diǎn)坐標(biāo)
float2 normalizedCoordinate = convert_float2(coordinate) * (float2)(2, 2);
//根據(jù)歸一化坐標(biāo)從原圖中讀取像素?cái)?shù)據(jù)
uint4 colour = read_imageui(in, sampler, convert_int2(normalizedCoordinate));
//將像素?cái)?shù)據(jù)寫(xiě)入目標(biāo)圖像
write_imageui(out, coordinate, colour);
}
實(shí)際比較，vx_not, vx_scale使用opencl, c_model實(shí)現(xiàn)時(shí)間對(duì)比：
opencl:
average time: 44099.857143 us
c_model:
average time: 68343.380952 us
7. vx debug print信息
程序中通過(guò)獲取VX_ZONE_MASK環(huán)境變量的值來(lái)設(shè)置Log級(jí)別，可以通過(guò)如下將所有級(jí)別信息都打開(kāi)：
export VX_ZONE_MASK=fffff
一共有如下幾個(gè)級(jí)別，每個(gè)級(jí)別占int的一個(gè)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實(shí)現(xiàn)：
https://github.com/GPUOpen-ProfessionalCompute-Libraries/amdovx-core
https://github.com/GPUOpen-ProfessionalCompute-Libraries/amdovx-modules
總結(jié)
- 上一篇： π 的定义（极限）
- 下一篇： 【hdu 6038】Function