简单的访问模式：

第一个也是最简单的合并案例可以通过任何支持CUDA的设备来实现：第k个线程访问缓存行中的第k个字。 并非所有线程都需要参与。例如，如果warp访问的线程相邻4字节字（例如，相邻浮点值），单个128B L1高速缓存线并因此单个合并事务将服务该存储器访问。 图3显示了这种模式。

如果线的某些字未被任何线程请求（例如，如果多个线程已经访问了相同的字或者某些线程没有参与该访问），则高速缓存线中的所有数据都被获取。 此外，如果在该段内已经对warp线程进行了访问，则只有一个128字节的L1事务将由具有计算能力2.x的设备执行。

顺序但未对齐的访问模式:如果顺序线程顺序但未与高速缓存线对齐，则会请求两个128字节的L1高速缓存，如图4所示。

对于非高速缓存事务（即绕过L1并仅使用L2高速缓存的事务），除32字节L2段的级别外，可以看到类似的效果。 在图5中，我们看到一个例子：使用与图4相同的访问模式，但现在禁用了L1缓存，因此现在需要五个32字节的L2段来满足请求。

通过CUDA Runtime API分配的内存（如通过cudaMalloc（））保证至少对齐至少256个字节。 因此，选择明显的线程块大小，例如翘曲大小的倍数（即，当前GPU上的32），便于通过与高速缓存线对齐的弯曲进行存储器访问。 （例如，如果线程块大小不是warp大小的倍数，请考虑第二个，第三个和后续线程块访问的内存地址会发生什么情况。）

错位访问的影响：使用简单的复制内核来探究未对齐的访问的后果是容易的和信息性的，例如复制内核中的一个，它说明未对齐的访问。说明错位访问的复制内核：

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

在说明未对齐访问的副本内核中，将数据从输入数组idata复制到输出数组，这两者都存在于全局内存中。 内核在主机代码的循环中执行，该代码将参数偏移量从0改为32.（图4和图4分别对应于缓存和非缓存内存访问情况下的错位）。副本的有效带宽 在NVIDIA Tesla M2090上具有各种偏移量（计算能力2.0，开启了ECC，因为它是默认设置）如图6所示。

对于NVIDIA Tesla M2090，全局内存访问没有偏移或偏移量为32个字的倍数，导致单个L1高速缓存线事务或4个L2高速缓存段负载（用于非L1高速缓存负载）。 实现的带宽约为130GB / s。 否则，每个warp都会加载两个L1高速缓存行（高速缓存模式）或者四个到五个L2高速缓存段（非高速缓存模式），从而实现了无偏移量的约四分之一的内存吞吐量。有趣的一点是，我们可能预期缓存情况比此示例的非缓存情况更差，因为缓存情况下的每个翘曲都会按需要获取两倍的字节数，而在非缓存情况下， 5/4根据需要获取的字节数是按每个warp获取的。 然而，在这个特殊的例子中，这种效果并不明显，因为相邻的经线会重复使用它们的邻居所获取的缓存线。 因此，虽然缓存负载的影响仍然很明显，但并不像我们预期的那么好。 如果相邻的经线没有表现出超高速缓存线的高度重用，情况会更加如此。