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教程