计算能力5.x
Architecture:
多处理器由以下部分组成:
- 用于算术运算的128个CUDA内核(请参见算术运算吞吐量的算术指令),
- 32个用于单精度浮点超越函数的特殊功能单元,
- 4个warp调度器
当多处理器被执行时,它首先在四个调度器中进行分配。 然后,在每个指令发布时间,每个调度程序为其分配的一个warp指令发出一条指令,如果有的话,该指令即可执行。
多处理器具有:
- 一个由所有功能单元共享的只读常量高速缓存,可加速驻留在设备内存中的常量内存空间的读取速度,
- 一个24 KB的统一L1 /纹理缓存,用于缓存全局内存中的读取,
- 计算能力5.0或96 KB共享内存的设备共享内存的64 KB共享内存用于计算能力的设备5.2。
纹理单元也使用统一的L1 /纹理缓存,实现纹理和表面存储器中提到的各种寻址模式和数据过滤。
还有一个由所有多处理器共享的L2缓存,用于缓存对本地或全局内存的访问,包括临时寄存器溢出。 应用程序可以通过检查l2CacheSize设备属性来查询L2缓存大小。
高速缓存行为(例如,读取是否高速缓存在统一的L1 /纹理高速缓存和L2中或L2中)可以使用对加载指令的修饰符以逐访问为基础部分配置。
全局内存:
全局内存访问总是缓存在L2中,而L2中的缓存与计算功能3.x的设备(参见全局内存)的行为方式相同。
只读内核整个生命周期的数据也可以通过使用__ldg()
函数(请参见只读数据高速缓存加载函数)读取,将其缓存在上一节所述的统一L1 /纹理高速缓存中。 当编译器检测到某些数据的只读条件满足时,它将使用__ldg()
来读取它。 编译器可能无法始终检测到某些数据的只读条件满足。 使用const和__restrict__
限定符标记用于加载此类数据的指针会增加编译器检测只读条件的可能性.
对于内核整个生命周期不是只读的数据不能缓存在计算能力为5.0的设备的统一L1 /纹理缓存中。 对于计算能力5.2的设备,默认情况下,它不会缓存在统一的L1 /纹理缓存中,但可以使用以下机制启用缓存:
- 使用PTX参考手册中所述的相应修改器,使用内联汇编执行读取;
- 使用-Xptxas -dlcm = ca编译标志进行编译,在这种情况下,所有读取都将被缓存,但使用内联汇编使用禁用缓存的修饰符执行的读取除外;
- 使用-Xptxas -fscm = ca编译标志进行编译,在这种情况下,所有读取都将被缓存,包括使用内联汇编执行的读取,而不管所使用的修改器如何。
当使用上面列出的三种机制启用缓存时,计算能力5.2的设备将全局内存读取缓存到统一的L1 /纹理缓存中,用于所有内核启动,除了内核启动之外,其内的线程块消耗太多的多处理器资源。 这些异常由分析器报告。
共享内存:
共享内存有32个组织,这样连续的32位字映射到连续的组。 每个bank每个时钟周期的带宽为32位。
对于一个warp的共享内存请求不会在访问同一个32位字内的任何地址的两个线程之间产生bank冲突(即使这两个地址落在同一个bank中):在这种情况下,对于读取访问, 被广播到请求线程和写入访问,每个地址仅由其中一个线程写入(该线程执行写入操作是未定义的)