《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.