CUDA实践指南（十五）

分段访问：
如上所示，在顺序访问未对齐的情况下，计算能力2.x设备的高速缓存有助于实现合理的性能。 然而，它可能与非单位跨步访问有所不同，并且这是处理多维数据或矩阵时经常发生的模式。 出于这个原因，确保所获取的每个缓存行中的数据尽可能多地被实际使用，这是这些设备上的存储器访问的性能优化的重要部分。
为了说明跨接访问对有效带宽的影响，请参阅内核中的内核strideCopy（），以说明非单元跨度数据副本，该数据副本用从idata到odata的线程之间的跨步元素跨度复制数据。
用于说明非单位步幅数据复制的内核:

__global__ void strideCopy(float *odata, float* idata, int stride)
{
    int xid = (blockIdx.x*blockDim.x + threadIdx.x)*stride;
    odata[xid] = idata[xid];

图7说明了这种情况; 在这种情况下，warp中的线程访问内存中的字，步长为2.此操作导致Tesla M2090上每个warp的两个L1高速缓存行（或非高速缓存模式下的八个L2高速缓存段）（计算能力2.0）

由于事务中的一半元素没有被使用并且代表浪费的带宽，所以2的跨度导致50％的加载/存储效率。 如图8所示，随着步幅的增加，有效带宽会减少，直到在32个线程中加载了32行高速缓存为止。

如图8所示，应尽可能避免非单位跨度全局内存访问。 一种方法是利用共享内存，这将在下一节讨论。
共享内存:
由于它是片上的，因此共享内存比本地和全局内存具有更高的带宽和更低的延迟 - 前提是线程之间不存在组间冲突，如下节所述:
共享内存和内存组:
为了实现并发访问的高内存带宽，共享内存被分成可以同时访问的大小相同的内存模块（库）。因此，跨越n个不同存储体的n个地址的任何存储器加载或存储可以被同时服务，产生的有效带宽是单个存储体带宽的n倍。
但是，如果一个内存请求的多个地址映射到同一个内存组，则访问将被序列化。硬件将存在bank冲突的内存请求拆分为许多单独的无冲突请求，并将有效带宽减少一个等于单独内存请求数量的因子。这里的一个例外是warp中的多个线程寻址相同的共享内存位置，导致广播。计算能力2.x及更高版本的设备具有多播共享存储器访问的附加能力（即，将相同值的副本发送到warp的多个线程）。
为了最大限度地减少银行冲突，了解内存地址如何映射到内存条以及如何最优地调度内存请求非常重要。
计算能力2.x:
在计算能力2.x的设备上，每个存储体每两个时钟周期的带宽为32位，连续的32位字被分配给连续的存储体。 warp的大小是32个线程，bank的数量也是32个，所以在warp中的任何线程之间都会发生bank冲突。
计算能力3.x:
在计算能力3.x的设备上，每个存储体在每个时钟周期（*）中具有64位的带宽。 有两种不同的银行模式：将连续的32位字（32位模式）或连续的64位字（64位模式）分配给连续的存储单元。 warp的大小是32个线程，bank的数量也是32个，所以在warp中的任何线程之间都会发生bank冲突。