CUDA学习——Chapter 2（4）内存空间布局对核函数性能的影响（1）

第二章

1.坐标映射

从前面的博文我们可以知道，global memory是可以划分成网格（一个程序对应一个网格），网格由块组成，块由线程组成。一个块内的线程可以相互访问，相互等待。

通过对前面并行向量加法的分析，我们知道，网格和块的大小会影响核函数的性能，这一篇博文将随着书本来探究如何组织网格和块从而获得更高效的性能。

还是以矩阵加法为例，在矩阵加法中，传统的是使用二维网格和二维块的布局来分配线程的归属，但这种分配方法不能获得最好的性能，我们可以自然地提出以下几种不同的分配方法：
1.由二维线程块构成的二维网格（就是矩阵里面套矩阵）
2.由一维线程块构成的一维网格（就是长条套长条）
3.由一维线程块构成的二维网格（就是矩阵里面套长条）

那么首先，我们要了解如何在各种奇奇怪怪的布局下得到线程的索引，该索引对于任意一个global memory分割方式来说，都是确定的。所以我更愿意将其称之为线程的坐标，但为了统一性，下文还是称作为索引。

对于一个矩阵来说，假定分割方式是二维线程块组成的二维网格，我们把线程当作矩阵中的一个元素。设某个线程在矩阵的坐标映射为(ix,iy)那么我们有

ix=threadIdx.x+blockIdx.x*blockDim.x
iy=threadIdx.y+blockIdx.y*blockDim.y

有点抽象，我们来具体地分析一下
这个式子其实写成
ix=blockIdx.x*blockDim.x+threadIdx.x
更好

首先，我们从大的单位：块，找起。线程位于第n个块的第i个线程，每个块有k个线程，那么它的坐标应该表示为（从1开始计数）：
x=n*k+i
这个还是很抽象，我们再具体一点

我们假设一个学校的一个年级有18个班，班号从1排到18，又假设每个班都有50人，学号从1排到50。现在要给全年级的学生分配一个唯一的年级识别号，那么16班32号的年级识别号就为：
16*50+32=832
也就是它在年级中的索引为832。
这样就好理解了吧？
y轴方向是一样的情况，就不再赘述了。

我们再往上看，刚才分析的是矩阵，现在我要将这个矩阵拍扁（就像将数字图像转化为一个unsigned char *一样），那么就把矩阵的每一行拿出来（从第2行开始），将它的头和上一行的尾相