CUDA学习--内存处理之共享内存（3）

1.共享内存

共享内存实际上是可受用户控制的一级缓存。每个SM中的一级缓存与共享内存共享一个64KB的内存段。

在实际中，共享内存的速度几乎在所有的GPU中都一致（大约为1.5TB/s的带宽），因为共享内存的速度受核时钟频率驱动。因此在任何显卡中，无论是否为高端显卡，除了使用寄存器外，还要更有效的使用共享内存。然而，GPU执行的是一种内存的加载-存储模型，即所有的操作都要在指令载入寄存器后才能执行。因此，加载数据到共享内存与加载数据到寄存器中不同，只有当数据重复利用、全局内存合并，或线程之间有共享数据时使用共享内存才更合适，否则，将数据直接从全局内存加载到寄存器性能会更好。

在编写代码时，可以将CUDA C的关键字shared添加到变量声明中，这将使这个变量驻留在共享内存中。在声明共享内存变量后，线程块中的每个线程都共享这块内存，使得一个线程块中的多个线程能够在计算上进行通信和协作。在世上没有免费的午餐，如果想要在线程之间进行通信，那么还需要一些操作来完成同步。例如，如果线程A将一个值写入到共享内存，并且我们希望线程B对这个值进行一些操作，那么只有当线程A的写入操作完成之后，线程B才能开始执行它的操作。
这里我们将调用__syncthread()函数来完成线程同步，这个函数的调用将确保线程块中的每个线程都执行完__syncthreads()前面的语句后，才会执行下一条语句。这样，可以确保对共享内存进行读取之前，想要写入的操作已经完成。但是，需要注意的是，CUDA架构将确保，除非线程块中的每个线程都执行__syncthreads()，否则没有任何线程能够执行__syncthreads()之后的指令。 当某些线程需要执行一条指令，而其他线程不需要执行时，这种情况就成为线程发散。也就是说，如果__syncthreads()位于发散分支，那么一些线程将永远无法执行__syncthreads()，那么这将使线程保持等待，等待，等待……所以，我们在使用__syncthreads()的时候要小心谨慎。

具体使用共享内存的用例请移至CUDA学习–矩阵乘法的并行运算。