关于《OPENCL异构并行计算》中卷积优化的分析

原创 2017年04月16日 17:13:39

《OPENCL异构并行计算》中讲了如何利用OPENCL进行卷积运算,并给出了使用局部存储器优化的例子,这里对其进行简单分析:

前面内容很容易懂,需要注意的是这里他计算卷积核的时候并没有提前把卷积核旋转180度,因此旋转操作需要自己来做。

对9.5节的局部存储器优化,他使用的原始图像,卷积核和输出的卷积MAP都是正方形的。

这里主要分析他是怎么初始化原始图像的,因为后面的计算过程一目了然。

T local l_pixels[(BS+filterSize-1)*(BS+filterSize-1)];

第一行就定义了局部存储器,这里需要说明一下,划分划分工作组的时候,工作组的大小是最终输出MAP的大小部分,举个例子,用5X5的卷积核来卷积28X28的图像,则最终输出MAP大小为24X24,然后进行局部优化的时候,我们可以令上面的BS=4,则是吧输出MAP,也就是24X24图像按照4X4一个工作组进行划分,而不是把原始图像按照4X4划分。这样每个工作组的大小为4X4,而每个工作组需要保存的原始图像大小就是(4+5-1)X(4+5-1),即把生成这4X4的输出MAP的那8X8的原始图像保存到局部存储器重以便访问。局部存储器的大小为(BS+filterSize-1)*(BS+filterSize-1),其中BS为输出卷积MAP的长或者宽(正方形都一样),filterSize为卷积核的边长,则BS+filterSize-1就是原始图像的边长(联系求输出MAP大小用的是原始图像边长-卷积核边长+1,这里正好反过来),因此这个局部数组用来存储的是整个原始图像的大小。

一定要注意的是这个函数的工作组的大小是输出卷积MAP的大小,而不是原始图像区域的大小,否则会出现计算错误。

接下来看索引:

int tidy = get_local_id(1);//工作项局部ID的y

int y = get_global_id(1);//工作项的全局ID的y

int tidx = get_local_id(0);// 工作项局部ID的x

int x = get_global_id(0);// 工作项的全局ID的x

int imageInSizeX = imageOutSizeX+filterSize-1;//输出全局MAP的边长

然后接下来就是把这一部分卷积MAP对应的原始图像放在局部存储器l_pixels中,例如我们的卷积核大小为5X5,原始图像大小为8X8,则输出图像大小为4X4,我们令BS=2,也就是把输出的4X4的MAP切分成4个2X2的MAP来计算,自然也需要4个工作组,以第二个工作组为例,原始图像为:

(0,0)

(1,0) 

(2,0) 

(3,0) 

(4,0) 

(5,0) 

(6,0) 

(7,0) 

(0,1) 

(1,1) 

(2,1) 

(3,1) 

(4,1) 

(5,1) 

(6,1) 

(7,1) 

(0,2) 

(1,2) 

(2,2) 

(3,2) 

(4,2) 

(5,2) 

(6,2) 

(7,2)

(0,3) 

(1,3) 

(2,3) 

(3,3) 

(4,3) 

(5,3) 

(6,3) 

(7,3) 

(0,4) 

(1,4) 

(2,4) 

(3,4) 

(4,4) 

(5,4) 

(6,4) 

(7,4) 

(0,5) 

(1,5) 

(2,5) 

(3,5) 

(4,5) 

(5,5) 

(6,5) 

(7,5) 

(0,6) 

(1,6) 

(2,6) 

(3,6) 

(4,6) 

(5,6) 

(6,6) 

(7,6) 

(0,7) 

(1,7) 

(2,7) 

(3,7) 

(4,7) 

(5,7) 

(6,7) 

(7,7) 

图中点(x,y)为坐标值,上图与大小为5X5的卷积核卷积,得到4X4大小的输出MAP:

(0,0)

(1,0) 

(2,0) 

(3,0) 

(0,1) 

(1,1) 

(2,1) 

(3,1) 

(0,2) 

(1,2) 

(2,2) 

(3,2) 

(0,3) 

(1,3) 

(2,3) 

(3,3) 

我们以第二个工作组为例,也就是图中涂色区域,看原始图像对应哪一部分:

对于(2,0)点,原始图像对应区域如下:

(0,0)

(1,0) 

(2,0) 

(3,0) 

(4,0) 

(5,0) 

(6,0) 

(7,0) 

(0,1) 

(1,1) 

(2,1) 

(3,1) 

(4,1) 

(5,1) 

(6,1) 

(7,1) 

(0,2) 

(1,2) 

(2,2) 

(3,2) 

(4,2) 

(5,2) 

(6,2) 

(7,2) 

(0,3) 

(1,3) 

(2,3) 

(3,3) 

(4,3) 

(5,3) 

(6,3) 

(7,3) 

(0,4) 

(1,4) 

(2,4) 

(3,4) 

(4,4) 

(5,4) 

(6,4) 

(7,4) 

(0,5) 

(1,5) 

(2,5) 

(3,5) 

(4,5) 

(5,5) 

(6,5) 

(7,5) 

(0,6) 

(1,6) 

(2,6) 

(3,6)

(4,6) 

(5,6) 

(6,6) 

(7,6) 

(0,7) 

(1,7) 

(2,7) 

(3,7) 

(4,7) 

(5,7) 

(6,7) 

(7,7) 

即从(2,0)开始的5X5大小的区域,照此我们可以画出

(2,0)

(3,0) 

(2,1) 

(3,1) 

区域对应原始图像的区域:

(0,0)

(1,0) 

(2,0) 

(3,0) 

(4,0) 

(5,0) 

(6,0) 

(7,0) 

(0,1) 

(1,1) 

(2,1) 

(3,1) 

(4,1) 

(5,1) 

(6,1) 

(7,1) 

(0,2) 

(1,2) 

(2,2) 

(3,2) 

(4,2) 

(5,2) 

(6,2) 

(7,2) 

(0,3) 

(1,3) 

(2,3) 

(3,3) 

(4,3) 

(5,3) 

(6,3) 

(7,3) 

(0,4) 

(1,4) 

(2,4) 

(3,4) 

(4,4) 

(5,4) 

(6,4) 

(7,4) 

(0,5) 

(1,5) 

(2,5) 

(3,5) 

(4,5) 

(5,5) 

(6,5) 

(7,5) 

(0,6) 

(1,6) 

(2,6) 

(3,6) 

(4,6) 

(5,6) 

(6,6) 

(7,6) 

(0,7) 

(1,7) 

(2,7) 

(3,7) 

(4,7) 

(5,7) 

(6,7) 

(7,7) 

即我们这一组的局部存储器l_pixels要存储的原始图像的部分区域就是上图的涂色区域。

接下来我们来思考一个问题,我们的工作组大小是2X2,而工作项需要的原始图像大小为6X6,那么怎么使用2X2的工作组来填充6X6的矩阵呢?

首先,我们知道我们的每一个工作项首先可以在初始化局部寄存器的时候,首先把他直接对应原始图像的那一块初始化了,如上图的第一个输出项(2,0)对应的原始图像卷积的第一个点是(2,0),则我们让局部存储器的(0,0)填充原始图像的(2,0)的内容,然后这样填充完了,我们的局部存储器就填充了原始图像的2X2的大小了,如下图所示:

A

B 

  

  

  

  

C 

D 

  

  

  

  

  

  

  

  

  

  

  

  

  

  

  

  

  

  

  

  

  

  

  

  

  

  

  

  

我们把四个工作项分别命名为ABCD,他负责填充的局部存储器的区域就用它的序号命名。然后他们首先可以填充自己对应的那个(tidx,tidy)位置的内容,自然是从原始图像的(x,y)处填充的。

然后我们得思考,这个时候四个工作项只填充了四个位置,还有36-4=32个位置没有填充,怎么办?

我们就得想着让一个工作项填充多个区域,填充的方式就是每个工作项负责填充他的X方向,Y方向,离他步数为BS的局部存储器对应的区域,这样每个工作项负责填充多个块,并且不会发生填充冲突,即我们每个工作项要负责的区域如下图:

A

B 

A 

B 

A 

B 

C 

D 

C 

D 

C 

D 

A 

B 

A 

B 

A 

B 

C 

D 

C 

D 

C 

D 

A 

B 

A 

B 

A 

B 

C 

D 

C 

D 

C 

D 

 

比如我们第一个工作项A负责填充的区域就是上图涂色部分,就是从他原始位置开始,沿各个方向,步数为2的填充对应的原始图像。

但是我们会发现这样写起代码来非常麻烦,我们想要的是工作项最多沿他的X方向,Y方向和对角线方向各填充一次就可以了,而不是填充这么多次,增加了代码的复杂度,则我们设计BS大小的时候就要考虑让局部存储器大小的一半小于等于BS了,即:

(BS+filterSize-1)/2<=BS,得BS>=filterSize-1,即我们需要让我们工作组的大小的边长大于等于卷积核的大小-1。

我们再举一个例子来说明这个问题,我们设原始图像大小为5X5,卷积核大小为2X2,则输出MAP大小为4X4,同样把BS设置为2,即把4X4大小划分成4个2X2大小的工作组,此时满足BS>=filterSize-1,即2>=2-1,而这个工作组所需要的局部内存大小为2+2-1=3,即3X3,则首先,我们让这四个工作项填充他们各自对应原始图像的位置:

A(0,0)

B(1,0)

(2,0)

C(0,1)

D(1,1)

(2,1)

(0,2)

(1,2)

(2,2)

然后,我们需要让他们填充他们沿三个方向平移BS所对应的位置的内容,由上图可知A沿水平方向+2为(2,0),那个位置可以去填充,因此A工作项负责填充(2,0),同理可以填充(0,2),(2,2);而B沿水平方向理论上应该填充(3,0),可是局部存储器没有(3,0),因此水平方向不需要B填充,而垂直方向B可以填充(1,2);同理C不能填充垂直方向,但可以填充水平方向的(2,1);最后,D工作项除了他本身之外,没有任何需要他填充的位置。即如果一个工作项可能在X方向上填充,则他的局部ID的x会小于filterSize-1;如果一个工作项可能在Y方向上填充,则他的局部ID的y会小于filterSize-1;如果一个工作项可能在对角线方向上填充,则他的局部ID的x和y都会小于filterSize-1。如图所示:

A(0,0)

B(1,0)

A(2,0)

C(0,1)

D(1,1)

C(2,1)

A(0,2)

B(1,2)

A(2,2)

这就是整个初始化矩阵所对应的算法,代码如下:

l_pixels[tidx+tidy*(BS+filterSize-1)] = imageIn[y*imageInSizeX+x];

    //right

    if(tidx < filterSize-1){

        l_pixels[tidx+BS+tidy*(BS+filterSize-1)] = imageIn[y*imageInSizeX+x+BS];

    }

    if(tidy < filterSize-1){

        l_pixels[tidx+(tidy+BS)*(BS+filterSize-1)] = imageIn[(y+BS)*imageInSizeX+x];

    }

    if(tidy < filterSize-1 && tidx < filterSize-1){

        l_pixels[tidx+BS+(tidy+BS)*(BS+filterSize-1)] = imageIn[(y+BS)*imageInSizeX+BS+x];

    }

最后不要忘了同步问题:

    barrier(CLK_LOCAL_MEM_FENCE);

即完成了局部存储器的初始化,然后接下来就是进行卷积操作然后放到相应的全局存储器里面,思路很简单这里就不做过多阐述了。

而对于9.6节的一个工作项对应多个输出,就是把相应的局部存储器的MAP多存几个,在两个方向上划分,一个工作项负责多个区域,理解起来很容易,不做过多叙述。

而注意划分BS大小的时候应被输出Map整除,否则会出现访问溢出问题,例如24X24的Map,应该设置BS大小为4而不是5,8也可以而不是9,总之要能被24整除。

This is OK.

相关文章推荐

OpenCL 优化后的卷积代码

自己修改过的 可以直接运行 卷积模板

OpenCL编程:图像卷积

图像卷积,就是对图像所有像素进行一些特定的运算处理。这里涉及两个问题,一是读取图片文件信息,二是作何种卷积运算。第一个问题可在《freeimage存取图片数据》里找到答案。第二个问题可以baidu卷积...
  • jaccen
  • jaccen
  • 2016年05月13日 17:35
  • 1349

使用OpenCL+OpenCV实现图像卷积(一)

基于VS2010,使用OpenCL+OpenCV实现图像卷积处理。

使用OpenCL+OpenCV实现图像卷积(二)

基于VS2010,使用OpenCL+OpenCV实现图像卷积处理。

MEC —— 优化内存与速度的卷积计算

本次介绍一种内存利用率高且速度较快的卷积计算方法。 来自ICML2017, 《MEC: Memory-efficient Convolution for Deep Neural Network》1. ...
  • shuzfan
  • shuzfan
  • 2017年08月20日 13:01
  • 1109

深度学习卷积算法的GPU加速实现方法

参考链接: http://blog.csdn.net/u010620604/article/details/52464529 http://blog.csdn.net/guoyilin...
  • mao_kun
  • mao_kun
  • 2017年02月05日 18:46
  • 3134

OpenCL 优化后的卷积代码

自己修改过的 可以直接运行 卷积模板

关于《OPENCL异构并行计算》中工作组间同步的分析

《OPENCL异构并行计算》中讲了如何利用使用同步锁的机制对各个工作组之间进行同步,这里对其进行简单分析...
内容举报
返回顶部
收藏助手
不良信息举报
您举报文章:关于《OPENCL异构并行计算》中卷积优化的分析
举报原因:
原因补充:

(最多只允许输入30个字)