openCL 低延迟,频繁调用实验方案

最新推荐文章于 2024-01-22 16:08:21 发布

zsyddl2

最新推荐文章于 2024-01-22 16:08:21 发布

阅读量1.1k

点赞数

分类专栏：图形图像深度学习 Opencl c++ 文章标签： OpenCL 优化低延迟

本文链接：https://blog.csdn.net/zsyddl2/article/details/81540696

版权

c++ 同时被 3 个专栏收录

20 篇文章 0 订阅

订阅专栏

图形图像

11 篇文章 0 订阅

订阅专栏

深度学习

8 篇文章 0 订阅

订阅专栏

openCl低延迟,频繁调用实验

1 opencl 可以利用gpu对计算进行加速.gpu对比cpu的特点,并行的核多,同步操作并行计算效率高(逻辑步调完全一致).串行逻辑比如cpu. 可以加速的算法:

1) 大量的数学运算.比如矩阵的加减乘除. 其中数据类型为float的单精度/半精度吞吐量最大.int不如float.

2) 可并行的排序/搜索.这是利用多核并行的特点.如果可以将算法按照2分方式计算.那算法在gpu的时间杂度可降低到log(n).典型的算法有双调排序.最小值,最大值,求和计算等

不过cpu与gpu的数据交互与通信,在时间上开销很大.并不会象cpu调用SIMD 指令可以直接加速.需要加速计算的数据场景有一定限制: 一次提交给gpu计算数据量相对大.比如gpu单线程执行指令至少在1000个以上.此为大致测试结果,基于opencl2.x. intel i7 7700 集成gpu.

对于pcie的独立显卡,未测试.按照其传输机制推测,只会比此延迟大.

2 gpu延迟实验分析

gpu与cpu通信最快的为集成核显, 因为可以共用内存空间,使用 opencl 2.x的 SVM共享机制,实现数据0拷贝. 也可以更方便的使用主机程序内存数据.但是目前7代i7仍然未支持主机虚拟地址共享.只能使用 clSVMAlloc方式.

在此方式下,(使用非SVM方式,延迟更高)

1)调用clEnqueueNDRangeKernel, opencl 核函数直接return,一次消耗0.050ms左右;

2)如果使用clFinish(queue);或者clWaitForEvents;消耗只时间为0.098ms左右.

由此看出,kernel的调用有0.01ms左右的基础开销.想要高频的少量数据计算,比如cpu计算不超过0.01ms的,将不会得到任何的加速效果.

SVM机制可以提供数据同步.所以考虑直接使用同步方式.通过变量判断kernel的执行结束.

1) 最简单的做法:使用 SVM内存,传入一个标志位,用于判断kernel结束.(此处为单线程方式,多线程需要使用一组标志变量)

int* pSync =

(int*)clSVMAlloc(

context, // the context where this memory is supposed to be used

CL_MEM_WRITE_ONLY | CL_MEM_SVM_FINE_GRAIN_BUFFER| CL_MEM_SVM_ATOMICS,

size * sizeof(int), // amount of memory to allocate (in bytes)

// 0 // alignment in bytes (0 means default)

// );

while (*pSync != FINISHED)

{//循环等kernel结束

}

kernel内部

kernel void svmbasic (global Element* elements, global float *dst, global atomic_int* pSync)

{

....

*pSync = FINISHED;

}

如此改造后,时间提升为 0.05xms基本省下clFinish时间.

2) 使用批量的方式提交数据.

这个方式不方便在于: 1)需要积累一定的数据才可以计算.2)如果后面的数据依赖于前面数据计算的结果,或者部分依赖.那不太能做到批量提交计算.

3) 尝试方法:利用类似cpu线程的等待方式,使gpu线程驻留,随时等待主机数据计算?看似一个很完美的方法,实际却行不通.

在kernel函数中使用while(1) 等待主机的同步通知,直接导致程序卡死,显卡卡死,整个操作系统卡死.只能对电脑强制重启.为什么会卡死呢?不详细分析了.大致原因是显卡的底层机制把,基于命令队列的方式.所以某一个kernel卡住,显卡整个卡死,导致系统卡死.可gpukernel又没有类似线程的sleep函数等方式,不能实现线程轮训调度.那么这条路死了?

4) 根据显卡的计算过程和测试结果分析,有一个折中方法,算是柳暗花明又一村. 思路是这样的:

条件:(1) 可以用SVM的内存同步方式,等待主机数据的多次提交.是多次提交,并不是永远等待

(2) 不可以卡死kernel函数,或者占用太久.

根据这个思路,可以在kernel使用循环,控制循环次数,比如10次.; 每次循环进行一次计算数据操作.该操作不可过度消耗时间.比如0.xms.. 控制整个循环执行结束不可超过10ms 左右.

kernel函数改造如下

int nLoop = 0;

int nLoopMax = 400;

//100 次会死机. 单次执行 20*10*1000;为极限

while ( nLoop < nLoopMax)

{//依照通讯执行10次. 时间很短. 0.52 ms 20: 0.97ms

int nSync = 0;

{//等待命令消息

nSync = atomic_load(pSync);

if (nSync >= 3)//表述主机通知结束任务

break;

if (nSync == 1)//表示主机通知数据已经更新,可以进行下一次计算

break;

} while (1);

if (nSync >= 3)

{//退出循环结束程序

break;

}

int i = 0;

while (i++ < 20)

{//

//fB = i+fC;

//fA += fB + fC;;

fA += fB + i + nLoop;

fA += fB + i + nLoop ;

fA += fB + i + nLoop;

}

nLoop++;

if (nLoop < nLoopMax)

{//设置为2 表示通知主机一次计算完成

atomic_store(

pSync, 2);

}

//通知主机, kernel退出

atomic_store(

pSync, 4);

主机程序大致为:

for (int i = 1; i <= nRunKernelTimes; i++)

{

pSync->store(1);

err = clSetKernelArgSVMPointer(kernel, 2, pSync);

SAMPLE_CHECK_ERRORS(err);

.......

int nLoop = 0;

int nRunKernelTimes = 0;

while ( nLoop++ <10000000

)

{//使用这种同步很快.几乎和不同步一样的时间

if (nSync == 2)

{//kernel'一次计算结束, 开始下次计算

nRunKernelTimes++;

pSync->store(1);

}

else if (nSync == 4)

{//kernel 退出

break;

}

..................

}

流程大致为:

主机->设置kernel数据,通知kernel可以执行,(设置标志为1).->主机等待gpu一次计算完成(等待gpu设置标志为2)执行其他指令,->再次通知gpu(设置为1) .

gpu->等待命令标志,如果为1表示可以执行->执行计算,计算结束设置标志为2;->进入下一次等待主机提交数据的命令

如果gpukernel计算结束.则设置命令标志为4.标志通知主机kernel结束.

如此改造后,可以实现细粒度的cpu gpu同步.需要注意几点:

1) 一次kernel开启执行,需要保证执行2000个命令以上.方可抵消启动kernel的损耗时间.

2) 一次kernel内部的循环执行的计算算法,最少应该在200个加法计算以上. 否则原子操作的内存同步过程,会消耗的时间比列太多.

经过测试,kernel循环400次,每次执行800个加法计算, 循环内平均一次执行为 0.0025ms

kernel的循环,以及每次循环内的算法指令个数,时间消耗关系

算法指令个数	Kernel中循环次数	一次kernel执行时间	一次算法时间	100个指令执行时间
200	100	0.21ms	0.002ms	0.001
200	200	0.38	0.0019	0.00095
200	400	0.52	0.0013	0.00065
400	400	0.68	0.0015	0.00037
800	400	1.03	0.0025	0.0003125
800	100	0.31	0.003	0.000375