全局内存:
计算能力3.x的设备的全局内存访问缓存在L2中,计算能力3.5或3.7的设备的全局内存访问也可缓存在上一节中描述的只读数据缓存中; 他们通常不会被L1缓存。 计算能力3.5的某些设备和计算能力3.7的设备允许通过-Xptxas -dlcm = ca选项将参与缓存L1中的全局内存访问缓存到nvcc。
高速缓存行是128字节,映射到设备内存中128字节对齐的段。 在L1和L2中缓存的内存访问都使用128字节的内存事务处理,而L2中缓存的内存访问仅使用32字节的内存事务处理。 例如,在分散存储器访问的情况下,L2中的缓存因此可以减少过取。
如果每个线程访问的单词大小大于4个字节,则warp的内存请求首先被分割为独立发出的单独的128字节内存请求:
- 两个内存请求,每个半个warp请求一个,如果大小为8个字节,
- 四个内存请求,每个四分之一warp一个,如果大小为16个字节。
然后将每个内存请求分解为独立发布的缓存行请求。 缓存行请求在缓存命中的情况下以L1或L2缓存的吞吐量服务,否则以设备内存的吞吐量服务。
请注意,线程可以按任何顺序访问任何单词,包括相同的单词。
如果一个warp执行的非原子指令写入全局内存中warp的多个线程中的同一个位置,则只有一个线程执行写操作,而哪个线程没有定义。
只读内核整个生命周期的数据也可以通过使用__ldg()
函数(请参阅只读数据高速缓存加载函数)读取,将其缓存在上一节中描述的只读数据高速缓存中。 当编译器检测到某些数据的只读条件满足时,它将使用__ldg()
来读取它。 编译器可能无法始终检测到某些数据的只读条件满足。 使用const和__restrict__
限定符标记用于加载此类数据的指针会增加编译器检测只读条件的可能性。
图16显示了全局内存访问和相应内存事务的一些示例。
共享内存:
共享内存有32个bank,有两种寻址模式,如下所述。
可以使用cudaDeviceGetSharedMemConfig()查询寻址模式,并使用cudaDeviceSetSharedMemConfig()进行设置(有关更多详细信息,请参阅参考手册)。 每个bank每个时钟周期的带宽为64位。
图17显示了一些跨接访问的例子。
图18显示了一些涉及广播机制的内存读取访问示例。
64位模式:
连续的64位字映射到连续的bank。
对于一个warp的共享内存请求不会在访问同一个64位字内任何子字的两个线程之间产生bank冲突(即使这两个子字的地址落在同一个bank中):在这种情况下 对于读取访问,将64位字广播给请求的线程,对于写入访问,每个子字只由其中一个线程写入(该线程执行写入操作是未定义的)。
32位模式:
连续的32位字映射到连续的bank。
对于一个warp的共享内存请求不会在访问同一32位字内的任何子字的两个线程之间或两个32位字中的索引i和j处于相同的64字对齐段内 (即第一个索引是64的倍数的段)并且使得j = i + 32(即使两个子字的地址落入同一个存储体):在这种情况下,对于读取访问,32 位字被广播给发出请求的线程和写入访问,每个子字只被其中一个线程写入(该线程执行写入操作是未定义的)