CUDA编程 基础与实践 学习笔记（七）

条过半…
7.1 全局内存的合并和非合并访问
从fermi开始，有了SM级别的L1 cache和设备层次的L2 cache. 有了访问全局内存的请求后，先L1，再L2，都不成功的话最后全局内存DRAM中取。一次数据传输在默认情况下是32B.
合并度：线程束请求的字节数和该请求导致的数据传输处理的字节数之比。可以视为一种资源利用率的表征。100%则为合并访问，否则为非合并访问。（为啥数据传输处理的字节数会多余请求的字节数，因为要像cpu的访存机制一样从全局取到cache中吗？答：这是基操，不在计算范围内。和访存地址有关。）

cudaMalloc()分配的内存首地址至少是256字节的整数倍。
（1）顺序的合并访问
（2）乱序的合并访问
（3）不对齐的非合并访问(地址错位，必须要多一次地址访问)

（4）跨越式非合并访问（数据是横着存的，结果竖着取数据）

（5）广播式非合并访问(例如，一次取了32字节的浮点数，但每次只用其中的一个浮点数，合并度为4/32=12.5%)

矩阵操作一般定义一个TILE来进行分块矩阵的操作。一个疑问：block size一般用多少合适？反正不能超过1024个线程。

书中矩阵转置的例子很有意思：

同为矩阵转置，同样都会包含一个合并操作，一个非合并操作，为什么第二个方法的速度比第一个方法快一倍？因为第二个方法虽然读操作是非合并的，但从后，编译器会判断一个全局内存变量是否在整个核函数范围内是可读的，会自动调用(帕斯卡(3.5)架构会自动调用， kepler,麦克斯韦框架需要手动调用)函数__ldg()读取全局内存（纹理内存/表面内存，和常量内存类似但容量更大）。所以，在不能满足读写操作都是合并操作的情况下，尽量满足写操作是合并操作，因为读操作有提速的办法。

这章内容比较简单，主要要注意顺序数据访问的方式，如果错位了或竖着读数据就可能要多访存一次，浪费时间。矩阵操作一般用TILE进行分块操作；读操作会使用到纹理内存，自动调用__ldg()函数来提速