本章学习一下在opencl中如何实现矩阵的转置，主要的技巧还是利用好local memory，防止bank conflit以及使得全局memory的读写尽量是合并(coalensing)读写。
我们的矩阵是一副二维灰度图像256*256，矩阵的转置也就是图像的转置。每个thread处理16(4*4)个pixel(uchar)，workgroup的size是(16,16)。
下面直接看shader代码：
uint wiWidth = get_global_size(0);
uint gix_t = get_group_id(0);
uint giy_t = get_group_id(1);
uint num_of_blocks_x = get_num_groups(0);
uint giy = gix_t;
uint gix = (gix_t+giy_t)%num_of_blocks_x;
uint lix = get_local_id(0);
uint liy = get_local_id(1);
uint blockSize = get_local_size(0);
uint ix = gix*blockSize + lix;
uint iy = giy*blockSize + liy;
int index_in = ix + (iy)*wiWidth*4;
// 通过合并读写把输入数据装入到lds中
int ind = liy*blockSize*4+lix;
block[ind] = input[index_in];
block[ind+blockSize] = input[index_in+wiWidth];
block[ind+blockSize*2] = input[index_in+wiWidth*2];
block[ind+blockSize*3] = input[index_in+wiWidth*3];
因为workgroup size是（16,16），所以lix,liy的取值范围都是0-15，下面我们通过图片看下，lix=0 liy=0,lix=1 liy=0时候，ind,以及index_in的值，从而得到输入图像数据如何映射到local memory中。
lix=0 liy=0
lix=1 liy=0
下面是影射关系，(0,0) thread处理的16个pixel用血红色表示，它们映射到lds的0bank和16bank，(1,0)thread处理的像素用绿色表示，它们映射到lds的bank1和bank17，有效的避免了bank conflit，而全局memory的访问不同thread对应连续的全局memory空间，可以实现合并读写，从而提高程序性能。
把转置的数据写到全局memory中的代码如下：
ix = giy*blockSize + lix;
iy = gix*blockSize + liy;
int index_out = ix + (iy)*wiWidth*4;
ind = lix*blockSize*4+liy;
uchar4 v0 = block[ind];
uchar4 v1 = block[ind+blockSize];
uchar4 v2 = block[ind+blockSize*2];
uchar4 v3 = block[ind+blockSize*3];
// 通过合并读写把lds中数据写回到全局memory中
output[index_out] = (uchar4)(v0.x, v1.x, v2.x, v3.x);
output[index_out+wiWidth] = (uchar4)(v0.y, v1.y, v2.y, v3.y);
output[index_out+wiWidth*2] = (uchar4)(v0.z, v1.z, v2.z, v3.z);
output[index_out+wiWidth*3] = (uchar4)(v0.w, v1.w, v2.w, v3.w);
对应copy关系图如下：
完整的代码请参考：
工程文件gclTutorial9
代码下载：
稍后提供
