《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 |
比如我们第一个工作项Ａ负责填充的区域就是上图涂色部分，就是从他原始位置开始，沿各个方向，步数为２的填充对应的原始图像。
但是我们会发现这样写起代码来非常麻烦，我们想要的是工作项最多沿他的Ｘ方向，Ｙ方向和对角线方向各填充一次就可以了，而不是填充这么多次，增加了代码的复杂度，则我们设计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.
306
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
