CUDA学习之第五章 共享内存和常量内存（二）

5.2 共享内存的数据布局

5.2.1 方形共享内存

声明二维共享内存变量（行主序存储）：

__shared__ int tile[N][N];

对于二维线程块，他有两种访问方式：

tile[threadIdx.y][threadIdx.x]
tile[threadIdx.x][threadIdx.y]

由于第一种访问方式邻近的线程会访问连续的地址，所以一个线程束内不会冲突存储体，所以他是比较好的方式。

5.2.1.1 行主序访问和列主序访问

原文举个例子来说明上面的第一个行主序时好的，给了一个线程块32*32的，一个二维共享内存tile[32][32]，在线程中给他赋值然后返给全局一维数组out。
通过共享内存读写和全局内存的写，在此中改变threadIdx.x和threadIdx.y 的顺序来比较时间，可以得知第一种更好。

5.2.1.2 按行主序写和列主序读

发现列主序读会存储体冲突。

5.2.1.3 动态共享内存

就是换成一维的了，结果和前面的一样，还是行主序的好。

5.2.1.4 填充静态声明的共享内存

将第二维度加了一个变成了tile[32][33]，可以发现两种方法都不会冲突了，但是要注意这是Fermi设备所以这样可以，对于Kepler设备中八位的内存会有差别，需要调整填充大小。

5.2.1.5 填充动态声明的共享内存

在计算一维索引时填充，需要注意在写到全局内存时要调整索引保证不会因为填充而写错。

5.2.1.6 方形共享内存内核性能的比较

得到结果是填充有效提高性能，因为他减少了存储体冲突，第二就是动态共享内存内核会增加少量的消耗。

5.2.2 矩形共享内存

这回是tile[Row][Col]，矩形的二维矩阵。和方形结果一样，后面不赘述了，代码中会有一些方便矩形的操作。