[参考译文] Linux/processor-SDK-AM57X：OpenCL+DSP 加速

您好！  

感谢您的及时回复。

我已阅读本章： software-dl.ti.com/.../Foundational_Components_OpenCV.html 。

然后我改进我的测试代码。  

我使用 CMEM 在 DSP 和 A15之间共享存储器、 并在 OpenCV 应用中编写 OpenCL 内核代码。

我将其编译到一个 exec bin：DSP_Accelerate_OpenCV_OpenCL。

然后、我找到了执行 OpenCL+DSP 加速时成本时间的根本案例。  

下面是我的改进代码：

主函数：

int main (int argc、char ** argv)
｛
     unsigned char * raW_YUV = NULL；
     int c =-1；
     struct timespec tp0、TP1、TP2、TP3、TP4、 TP5；

void *cmem_src  =__malloc_DDR (camera_height * camera_width*2)；
void *cmem_dest =__malloc_dr (camera_height * camera_width*3)；
void *cmem_out  =__malloc_DDR (camera_height * camera_width*3)；

cameraOpenDevice (1)；

while (1)｛
     Clock_gettime (clock_monotonic、&tp0)；
     时间=(double) GetTickCount()；

     DqBuffer (RAW_YUV)；
     QBuffer()；

     MAT MAT_YUV (camera_height、camera_width、CV_8UC2)；
     memcpy (MAT_YUV.data、RAW_YUV、camera_height * camera_width*2)；

     Clock_gettime (clock_monotonic、&TP1)；

     MAT MAT_BGR (camera_height、camera_width、CV_8UC3、cmem_out)；
     ProcRawClcmem (MAT_YUV、"cvhighlighted.cl"、cmem_src、cmem_dest、MAT_BGR)；

     Clock_gettime (clock_monotonic、&TP2)；

     imshow ("frame"、MAT_BGR)；
     Clock_gettime (clock_monotonic、&TP3)；

     printf ("获取 RAW_YUV Mat tDIFF=%LF ms \n"、tDIFF_calc (tp0、TP1))；
     printf ("ProRawClcmem tdiff=%LF ms \n"、tDIFF_calc (TP1、TP2)；
     printf ("imshow tdiff=%LF ms \n"、tDIFF_calc (TP2、TP3))；

     C = waitKey (1)；
     if (c =27 || c ='q'|| c ='q')
     中断；
｝

     cameraCloseDevice()；

     _free_DDR (cmem_src)；
     _free_DDR (cmem_dest)；
     _free_DDR (cmem_out)；
     返回0；
｝

void ProcRawClcmem (Mat &mat_in、const std：：：string &kernel_name、void *cmem_src、void *cmem_dest、Mat &mat_out)
｛
     Int 错误；
     struct timespec tp0、TP1、TP2、TP3、TP4、 TP5、TP6、TP7、TP8、TP9、 tp10；

     MAT MAT_src (camera_height、camera_width、CV_8UC2、cmem_src)；
     MAT MAT_DEST (camera_height、camera_width、CV_8UC3、cmem_dest)；

     Clock_gettime (clock_monotonic、&tp0)；
     MAT_IN.CopyTo (MAT_src)；
     Clock_gettime (clock_monotonic、&TP1)；

     /R/W
     上下文上下文(CL_DEVICE_TYPE_Accelerator)；
     std：：向量 设备= context.getinfo ()；

     ifstream IFS (kernel_name)；
     如果(IFS.fail()){
          COUT <<"OpenCL 文件不可用..." <endl;
     ｝
     std：：string kernelStr (istreambuf_iterator (IFS))、istreambuf_iterator ())；

程序：：source (1、std：：：make_pair (kernelStr.c_str()、kernelStr.length())；
方案方案=方案(背景、来源)；
program.build(设备)；

内核内核(程序、"YUV2RGB_CMEM")；
CommandQueue 队列(上下文、设备[0]、CL_Queue_profiling_enable)；
Clock_gettime (clock_monotonic、&TP2)；

cl_mem bufSrc = clCreateBuffer (context()、CL_MEM_READ_ONLY |CL_MEM_USE_HOST_PTR、BUF_SRC、cmem_src、&err)；
cl_mem bufDst = clCreateBuffer (context()、cl_MEM_write_only|cl_MEM_use_host_ptr、BUF_dest、cmem_dest、&err)；

cl_event ev0；
ERR = clEnqueueWriteBuffer (queue()、bufSrc、CL_true、0、BUF_SRC、 MAT_src.data、0、NULL、&ev0)；
clWaitForEvents (1、&ev0)；
Clock_gettime (clock_monotonic、&TP3)；

ERR = clSetKernelArg (kernel()、0、sizeof (cl_mem)、&bufSrc)；
ERR = clSetKernelArg (kernel()、1、sizeof (cl_mem)、&bufDst)；

CL_EV1；
std：：size_t GlobalWokrSize[2]=｛(std：：size_t) camera_width、(std：size_t) camera_height｝；
ERR = clEnqueueNDRangeKernel (queue()、kernel()、2、NULL、GlobalWokrSize、 空、0、空、&EV1)；
clWaitForEvents (1、&EV1)；
Clock_gettime (clock_monotonic、&TP4)；

CL_EV2；
ERR = clEnqueueReadBuffer (queue()、bufDst、CL_true、0、BUF_dest、 MAT_DEST.DATA、0、NULL、&EV2)；
clWaitForEvents (1、&EV2)；

Clock_gettime (clock_monotonic、&TP5)；

ERR = clReleaseMemObject (bufSrc)；
ERR = clReleaseMemObject (bufDst)；
Clock_gettime (clock_monotonic、&TP6)；

MAT_DEST.CopyTo (MAT_OUT)；
Clock_gettime (clock_monotonic、&TP7)；

     /R/W
     printf ("Mat2Mat tdiff=%LF ms \n"、tDIFF_calc (tp0、TP1))；                 //0.472872ms
     printf ("构建内核 tDIFF=%LF ms \n"、tDIFF_calc (TP1、TP2)；                //3.456178毫秒
     printf ("clEnqueueWriteBuffer tDIFF=%LF ms \n"、tDIFF_calc (TP2、TP3))；       //0.608536毫秒
     printf ("clEnqueueNDRangeKernel tDIFF=%LF ms \n"、tDIFF_calc (TP3、TP4))；  //97.881413ms
     printf ("clEnqueueReadBuffer tDIFF=%LF ms \n"、tDIFF_calc (TP4、TP5))；       //0.871893ms
     printf ("clReleaseMemObject tDIFF=%LF ms \n"、tDIFF_calc (TP5、TP6))；       //0.011387毫秒
     printf ("Mat2Mat tdiff=%LF ms \n"、tDIFF_calc (TP6、TP7))；                 //1.283602ms
｝

cvhighlighted.cl：

__kernel void YUV2RGB_CMEM (__global const uchar* srcptr、__global uchar* dstptr)
｛
     int use_optimed_load = 0；

int src_offset = 0；
int dst_offset = 0；
int src_step = 640*2；
int dst_step = 640*3；
int rows = 480；
int cols = 640；

int x = get_global_id (0)；
int y = get_global_id (1)；

如果(x < cols /2)
｛
__global const uchar* src = srcptr + mad24 (y、src_step、(x << 2)+ src_offset)；
__global uchar* dst = dstptr + mad24 (y、dst_step、mad24 (x << 1、3、dst_offset))；

pragma 展开
对于(int cy = 0；cy < 1；+cy)
｛
如果(y < rows)｛
__constant float* coeffs = c_YUV2RGBCoeffs_420；
int load_src =*(__global int*) src)；
float VEC_src[4]=｛load_src & 0xff、(load_src >> 8)& 0xff、(load_src >> 16)& 0xff、(load_src >> 24)& 0xff｝；
float U = vec_src[1]- 128；
浮点 V = VEC_src[3]- 128；
浮点 Y00 = max (0.f、vec_src[0]- 16.f)* coeffs[0]；
float Y01 = max (0.f、vec_src[2]- 16.f)* coeffs[0]；

float RUV =((coeffs[4]*V)+0.5f)；
float guv =(coeffs[3]*V)+(coeffs[2]*U)+0.5f);
float BUV =((coeffs[1]*U)+0.5f)；

     dst[2]= convert_uchar_sat (Y00 + RUV)；
     dst[1]= convert_uchar_sat (y00 + guv)；
     dst[0]= convert_uchar_sat (Y00 + BUV)；
     dst[5]= convert_uchar_sat (Y01 + RUV)；
     dst[4]= convert_uchar_sat (y01 + guv)；
     dst[3]= convert_uchar_sat (Y01 + BUV)；
     ｝
     ++y；
     src += src_step；
     dst += dst_step；
     ｝
     ｝
｝

结果：

root@ok5718-idk：/home/forlinx/qt
root@ok5718-idk：/home/forlinx/qt
root@ok5718-idk：/home/forlinx/qt
root@ok5718-idk：/home/forlinx/qt
root@ok5718-idk：/home/forlinx/qt on_dsp_accelerate.sh

获取 RAW_YUV Mat tdiff = 4.396880ms
ProcRawClmem tdiff = 106.702501ms
imshow tdiff=1.879939ms

Mat2Mat tdiff = 0.455792ms
生成内核 tdiff=8.894614 ms
clEnqueueWriteBuffer tdiff = 1.317600 ms
clEnqueueNDRangeKernel tDIFF=98.243997ms
clEnqueueReadBuffer tDIFF = 1.948584ms
clReleaseMemObject tDIFF=0.004392 ms
Mat2Mat tdiff = 1.185027 ms

root@ok5718-idk：/home/forlinx/qt
root@ok5718-idk：/home/forlinx/qt
root@ok5718-idk：/home/forlinx/qt
root@ok5718-idk：/home/forlinx/qt
root@ok5718-idk：/home/forlinx/qt cat on_dsp_accelerate.sh
导出 TI_OCL_cache_kernels=Y
导出 OpenCV_OpenCL_DEVICE_='TI AM57：加速器：TI 多核 C66 DSP'
回显"OpenCL 打开"
/DSP_Accelerate_OpenCV_OpenCL

如您所见，函数 clEnqueueNDRangeKernel()在 DSP 工作时花费的时间最多。

我的问题是：

是否有任何方法 可以提高 clEnqueueNDRangeKernel 函数的效率？

谢谢。

您好！
clEnqueueNDRangeKernel()调用可能包含内核加载时间。
o请提升程序构建和加载时间测量值。 可以通过在同一程序中调用一个 null ()内核函数来加载程序。 有关详细信息，请参阅《OpenCL 用户指南》中的“null”示例
downloads.ti.com/.../overview.html
o检查编译的内核是否通过管道传递软件。 使用“-k”保留生成的汇编文件。
o如果需要进一步的内核优化，请尝试使用 EDMA/OCopies 进行双缓冲。 有关详细信息，请参阅“conv1d”示例和 TI 在线文档。

您的 Processor SDK 已旧、如果可以、请迁移到最新的5.3版本。

雷克斯

您好！
我听从您的建议、下面是结果：

OpenCL" null "示例

root@ok5718-idk：/home/forlinx/qt
root@ok5718-idk：/home/forlinx/qt
root@ok5718-idk：/home/forlinx/qt #./null
OpenCL 运行时会将器件程序延迟加载到上
首先从程序中排队内核、然后是经过的时间
整体而言、第一个排队的时间将更长、以应对
加载内核。 此外、内核的后续 enqueue 将会
还可能受益于器件上的热缓存。

已用(带负载)：1008 usecs
已用(无负载)：286 usecs
空内核执行：要提交的队列：2us
空 Kernel Exec：提交至 Start：30us
空内核执行：开始到结束：214us

已用(无负载)：279 usecs
空内核执行：要提交的队列：2us
空 Kernel Exec：提交至 Start：31us
空内核执行：开始到结束：210us

已用(无负载)：269 μ s
空内核执行：要提交的队列：2us
空 Kernel Exec：提交至 Start：30us
空 Kernel Exec：开始到结束：202us

已用(无负载)：280 μ s
空内核执行：要提交的队列：1us
空 Kernel Exec：提交至 Start：29us
空 Kernel Exec：开始到结束：215us

已用(无负载)：277 μ s
空内核执行：要提交的队列：1us
空 Kernel Exec：提交至 Start：28us
空内核执行：开始到结束：214us

完成！



2.使用"-k"保留生成的汇编文件。

我使用以下命令编译 OpenCL 内核：clocl -t -k cvhighlights.cl

我修改了我的代码：
#include "cvhighlighteds.dsp_h"
程序：二进制文件(1、std：：：make_pair (cvhighlighted_dsp_bin、sizeof (cvhighlighted_dsp_bin))；
程序 程序=程序(上下文、器件、二进制)；
program.build(设备)；

然后我在 shell 中执行它、结果如下：

构建内核 tdiff=1.538664ms
clEnqueueWriteBuffer tdiff=0.62399ms
clEnqueueNDRangeKernel tDIFF=97.438472ms
clEnqueueReadBuffer tdiff=0.905565ms
clReleaseMemObject tdiff=0.007808ms

clEnqueueNDRangeKernel()成本时间没有变化。
因此根本原因不是内核加载时间。


我发现真正的成本时间是内核函数__kernel void YUV2RGB_CMEM (__global const uchar* srcptr、__global uchar* dstptr)。

此函数来自 OpenCV 文件 cvhighlighted.cl 中的 YUV2RGB_422()函数：
__kernel void YUV2RGB_422 (__global const uchar* srcptr、int src_step、int src_offset、
_global uchar* dstptr、int dst_step、int dst_offset、
int 行、int cols)

我将其复制到了内核函数 YUV2RGB_CMEM()中。
它在使用浮点运算的 DSP 上执行真的需要太多的时间。

您好！
我发现了有关 OpenCL 内核函数 YUV2RGB_CMEM 的新问题：

原始 YUV2RGB_CMEM()函数为：

__kernel void YUV2RGB/ 422_CMEM (__global const uchar* srcptr、__global uchar* dstptr)
｛
int x = get_global_id (0)；
int y = get_global_id (1)；

如果(x < 320)｛
__global const uchar* src = srcptr +(y*1280+(x <<2)；
__global uchar* dst=dstptr +(y*1920 +((x << 1)*3))；

如果(y < 480)｛
__constant float* coeffs = c_YUV2RGBCoeffs_420；
int load_src =*(__global int*) src)；
float VEC_src[4]=｛load_src & 0xff、(load_src >> 8)& 0xff、(load_src >> 16)& 0xff、(load_src >> 24)& 0xff｝；
float U = vec_src[1]- 128；
浮点 V = VEC_src[3]- 128；
浮点 Y00 = max (0.f、vec_src[0]- 16.f)* coeffs[0]；
float Y01 = max (0.f、vec_src[2]- 16.f)* coeffs[0]；
float RUV =((coeffs[4]*V)+0.5f)；
float guv =(coeffs[3]*V)+(coeffs[2]*U)+0.5f);
float BUV =((coeffs[1]*U)+0.5f)；

////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
dst[2]= convert_uchar_sat (Y00 + RUV)；
dst[1]= convert_uchar_sat (y00 + guv)；
dst[0]= convert_uchar_sat (Y00 + BUV)；
dst[5]= convert_uchar_sat (Y01 + RUV)；
dst[4]= convert_uchar_sat (y01 + guv)；
dst[3]= convert_uchar_sat (Y01 + BUV)；
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
｝
｝
｝

使用此 OpenCL 内核，clEnqueueNDRangeKernel()的成本为：
clEnqueueNDRangeKernel tdiff = 98.937283ms

我在 shell 中使用 Linux top 命令、发现 cpu%为31.0。

然后我按照以下方式修改了 YUV2RGB/ CMEM()：

__kernel void YUV2RGB/ 422_CMEM (__global const uchar* srcptr、__global uchar* dstptr)
｛
int x = get_global_id (0)；
int y = get_global_id (1)；

如果(x < 320)｛
__global const uchar* src = srcptr +(y*1280+(x <<2)；
__global uchar* dst=dstptr +(y*1920 +((x << 1)*3))；

如果(y < 480)｛
__constant float* coeffs = c_YUV2RGBCoeffs_420；
int load_src =*(__global int*) src)；
float VEC_src[4]=｛load_src & 0xff、(load_src >> 8)& 0xff、(load_src >> 16)& 0xff、(load_src >> 24)& 0xff｝；
float U = vec_src[1]- 128；
浮点 V = VEC_src[3]- 128；
浮点 Y00 = max (0.f、vec_src[0]- 16.f)* coeffs[0]；
float Y01 = max (0.f、vec_src[2]- 16.f)* coeffs[0]；
float RUV =((coeffs[4]*V)+0.5f)；
float guv =(coeffs[3]*V)+(coeffs[2]*U)+0.5f);
float BUV =((coeffs[1]*U)+0.5f)；

//////////////////////////////////////////////////////////////////////////////////////////////////////////////////// 修改了 START
int Y1、Y2、Y3、Y4、Y5、Y6；
Y1 =(int)(Y00 + RUV)；
y2 =(int)(y00 + guv)；
Y3 =(int)(Y00 + BUV)；
Y4 =(int)(Y01 + RUV)；
y5 =(int)(y01 + guv)；
Y6 =(int)(Y01 + BUV)；

dst[2]=(uchar)(y1 < 0？ 0：(Y1 > 255？ 255：y1)；
dst[1]=(uchar)(Y2 < 0？ 0：(Y2 > 255？ 255：y2)；
dst[0]=(uchar)(Y3 < 0？ 0：(Y3 > 255？ 255：Y3))；
dst[5]=(uchar)(Y4 < 0？ 0：(Y4 > 255？ 255：Y4))；
dst[4]=(uchar)(Y5 < 0？ 0：(Y5 > 255？ 255：Y5))；
dst[3]=(uchar)(Y6 < 0？ 0：(Y6 > 255？ 255：Y6))；
//////////////////////////////////////////////////////////////////////////////////////////////////////////////////// 修改后的结束
｝
｝
｝

clEnqueueNDRangeKernel()成本时间为：
clEnqueueNDRangeKernel tdiff = 18.490971ms。

然后、我在 shell 中使用 top 命令、发现 cpu%为58.4。

虽然其成本时间大大减少、但 DSP 加速效应并不明显。

我只需将函数 convert_uchar_sat()作为其在/usr/share/ti/OpenCL/DSP.h 中的定义进行扩展、但它们的性能差异很大。

您能不能向我解释一下它的根本原因、非常感谢。