计算能力为3.x
Architecture:
多处理器由以下部分组成:
- 用于算术运算的192个CUDA内核(请参阅算术运算吞吐量的算术指令),
- 32个特殊功能单元用于单精度浮点超越函数
- 4个warp调度器。
当多处理器被执行时,它首先在四个调度器中进行分配。 然后,在每个指令发布时间,每个调度程序为其准备执行的其中一个分配的warp发出两条独立指令(如果有的话)
多处理器具有只读常量高速缓存,该高速缓存由所有功能单元共享,并加速驻留在设备内存中的常量内存空间的读取速度。
每个多处理器都有一个L1高速缓存,所有多处理器都有一个L2高速缓存。 一级缓存用于缓存对本地内存的访问,包括临时注册溢出。 二级缓存用于缓存对本地和全局内存的访问。 高速缓存行为(例如,读取是否缓存在L1和L2中,还是L2中)可以使用对加载或存储指令的修饰符在每次访问的基础上部分配置。 计算能力3.5的一些设备和计算能力3.7的设备允许通过编译器选项在L1和L2中选择缓存全局存储器。
// Device code
__global__ void MyKernel()
{
...
}
// Host code
// Runtime API
// cudaFuncCachePreferShared: shared memory is 48 KB
// cudaFuncCachePreferEqual: shared memory is 32 KB
// cudaFuncCachePreferL1: shared memory is 16 KB
// cudaFuncCachePreferNone: no preference
cudaFuncSetCacheConfig(MyKernel, cudaFuncCachePreferShared)
默认的缓存配置是“prefer none”,意思是“无偏好”。 如果内核被配置为没有首选项,那么它将默认为使用cudaDeviceSetCacheConfig()/ cuCtxSetCacheCo fig()(详见参考手册)设置的当前线程/上下文的首选项。 如果当前线程/上下文也没有首选项(这又是默认设置),那么任何内核最近使用的缓存配置将是使用的那个缓存配置,除非启动内核需要不同的缓存配置 例如,由于共享内存要求)。 初始配置是48 KB的共享内存和16 KB的L1缓存。
计算能力设备3.7为每个上述配置添加了额外的64 KB共享内存,分别为每个多处理器产生112 KB,96 KB和80 KB的共享内存。 但是,每个线程块的最大共享内存仍为48 KB。
应用程序可以通过检查l2CacheSize设备属性来查询L2缓存大小。 最大二级缓存大小为1.5 MB。
每个多处理器都有一个48 KB的只读数据高速缓存,用于加速从设备内存读取数据。 它直接访问这个缓存(对于计算能力3.5或3.7的设备),或者通过一个纹理单元来实现纹理和表面存储器中提到的各种寻址模式和数据过滤。 当通过纹理单元访问时,只读数据高速缓存也被称为纹理高速缓存。