关于《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编程:图像卷积

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

OpenCL与CNN篇四:CNN从入门到使用

记录我从零到实现一个具体CNN网络中最有用的知识干货。 以细节为切入点,分享我对CNN网络的简洁。 本文致力于让你一篇文章理解CNN的具体实现与训练方法。 涉及理论不一一追述背景,主要讲解其如何应用。...
  • ShiAokai
  • ShiAokai
  • 2017年06月10日 01:08
  • 889

OpenCL 优化后的卷积代码

自己修改过的 可以直接运行 卷积模板
  • chao56789
  • chao56789
  • 2016年01月07日 19:49
  • 1103

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

基于VS2010,使用OpenCL+OpenCV实现图像旋转功能。
  • icamera0
  • icamera0
  • 2017年06月08日 08:05
  • 397

CUDA和OpenCL异同点比较

CUDA和OpenCL异同点比较 一、概述    对CUDA和opencl有一定的编程经验,但是细心的人可以发现,OPENCL就是仿照CUDA来做的。既然两个GPU的编程框架如此相像,那么他们究竟有什...
  • zhouxuguang236
  • zhouxuguang236
  • 2014年12月16日 21:25
  • 4416

openCL 优化

openCL 优化 1 LocalMemory 避免 bank conflict 每个SIMD 引擎32KB 的 LDS 被分为 32 个 Bank 每个 Bank 的带宽是4byte的数...
  • chao56789
  • chao56789
  • 2015年02月10日 17:06
  • 843

OpenCL性能对比测试案例,合理使用clfinish()函数,充分发挥指令队列(commandqueue)特性。

用的公司的电脑,配置如下: CPU: AMD Athlon X4 830 (3.0GHz 四核) 内存: 8GB GPU: nVIDIA GT710 (0.954GHz 192cores 1CU...
  • jinxiu0406
  • jinxiu0406
  • 2017年09月20日 10:19
  • 486

opencl:提高kernel有效并发约束(Effective concurrency constraint)的一点心得

下面是一个简单的kernel函数,从integ_count_mat矩阵中用vload8函数取出A1,A2,A3,A4四个向量执行A4+A1-A2-A3,结果存入density_mat,代码中只用到了一...
  • 10km
  • 10km
  • 2016年05月09日 14:51
  • 1574

opencl:慎用-cl-opt-disable选项编译kernel(可能会导致一些无法解释的问题)

在编译opencl kernel代码时,有一个编译选项-cl-opt-disable。根据opencl 官网的原文描述,使用这个选项可以关闭所有的代码优化,便于调试程序。(默认情况下,编译优化选项是打...
  • 10km
  • 10km
  • 2016年04月21日 15:00
  • 1284

C++实战之OpenCL矩阵相乘优化(二)

前言上一篇文章,分析了简单的矩阵相乘在opencl里面的优化kernel代码,每个work-item只负责计算结果矩阵的一个元素。下一步准备每次计算出结果矩阵的块元素,看看计算时间是如何。这个矩阵系列...
  • w401229755
  • w401229755
  • 2017年11月21日 14:51
  • 1022
内容举报
返回顶部
收藏助手
不良信息举报
您举报文章:关于《OPENCL异构并行计算》中卷积优化的分析
举报原因:
原因补充:

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