1) SM包括: 128个用于算术运算的CUDA核(参见算术运算吞吐量的算术指令), 32个用于单精度浮点超越函数的特殊函数单元, 4个曲速调度器。 2) 当SM被赋予要执行的扭曲时,它首先在四个调度器之间分配它们。然后,在每个指令发布时间,每个调度器为其分配的准备执行的扭曲之一(如果有的话)发布一条指令。 SM具有: 由所有功能单元共享并且加速从驻留在设备存储器中的恒定存储器空间的读取的只读恒定高速缓存, 24KB的统一L1/纹理高速缓存,用于高速缓存来自全局存储器的读取, 用于计算能力为5.0的设备的64KB共享存储器或用于计算能力5.2的设备的96KB共享存储器。 3) 统一的L1/纹理缓存也由纹理单元使用,该纹理单元实现纹理和表面存储器中提到的各种寻址模式和数据过滤。 还有一个由所有SM共享的L2缓存,用于缓存对本地或全局内存的访问,包括临时寄存器溢出。应用程序可以通过检查l2CacheSize设备属性来查询二级缓存大小(请参阅设备枚举)。 缓存行为(例如,读取是缓存在统一的L1/纹理缓存和L2中还是仅缓存在L2中)可以使用加载指令的修饰符在每次访问的基础上进行部分配置。 4) 全局内存访问总是缓存在L2中。
在内核的整个生命周期内只读的数据也可以通过使用__ldg()函数读取,缓存在上一节中描述的统一L1/纹理缓存中(请参阅只读数据缓存加载函数)。当编译器检测到某些数据满足只读条件时,它将使用__ldg()读取该数据。编译器可能无法始终检测到某些信息满足只读条件。用const和__restrict__限定符标记用于加载此类数据的指针会增加编译器检测到只读条件的可能性。
对于计算能力为5.0的设备,在内核的整个生命周期内不是只读的数据不能缓存在统一的L1/纹理缓存中。对于计算能力为5.2的设备,默认情况下,它不缓存在统一的L1/纹理缓存中,但可以使用以下机制启用缓存:
按照PTX参考手册中的描述,使用带有适当修改器的内联组件进行读取;
使用-Xptxas-dlcm=ca编译标志进行编译,在这种情况下,所有读取都被缓存,但使用带禁用缓存的修饰符的内联汇编执行的读取除外;
使用-Xptxas-fscm=ca编译标志进行编译,在这种情况下,所有读取都被缓存,包括使用内联汇编执行的读取,而不管使用了什么修饰符。
当使用上面列出的三种机制之一启用缓存时,具有计算能力5.2的设备将在统一L1/纹理缓存中缓存所有内核启动的全局内存读取,除了线程块消耗过多SM寄存器文件的内核启动之外。探查器会报告这些异常。 5) Shared Memory共享存储器具有32个组,这些组被组织为使得连续的32位字映射到连续的组。每个存储体具有每个时钟周期32比特的带宽。
对扭曲的共享内存请求不会在访问同一32位字内任何地址的两个线程之间产生组冲突(即使这两个地址位于同一组中)。在这种情况下,对于读取访问,字被广播到请求线程,而对于写入访问,每个地址仅由其中一个线程写入(哪个线程执行写入是未定义的)。 ---) CUDA为每个线程块保留1 KB的共享内存。因此,A100 GPU使单线程块能够寻址高达163KB的共享存储器,并且具有计算能力8.6的GPU可以在单线程块中寻址高达99KB的共享内存。为了保持体系结构兼容性,静态共享内存分配仍限制在48 KB以内,并且还需要显式选择加入以启用高于此限制的动态分配。有关详细信息,请参阅《CUDA C++编程指南》。 |
说点什么...