1） A Streaming Multiprocessor (SM) consists of:
2） SM在其调度器之间静态地分配其扭曲。然后，在每个指令发布时间，每个调度器为其分配的准备执行的扭曲之一（如果有的话）发布一条指令。
SM具有：
由所有功能单元共享并且加速从驻留在设备存储器中的恒定存储器空间的读取的只读恒定高速缓存，
一个统一的数据高速缓存和共享存储器，其总大小为192KB，用于计算能力8.0和8.7的设备（1.5倍Volta的128KB容量）。
共享内存是从统一的数据缓存中分割出来的，可以配置为各种大小（请参阅共享内存部分）。剩余的数据缓存用作L1缓存，也由纹理单元使用，该纹理单元实现纹理和表面存储器中提到的各种寻址和数据过滤模式。 3) 4) Shared Memory与Volta体系结构类似，为共享内存保留的统一数据缓存的数量可以在每个内核的基础上进行配置。对于NVIDIA Ampere GPU架构，对于计算能力为8.0和8.7的设备，统一数据缓存的大小分别为192KB和128KB。对于计算能力8.0和8.7的设备，共享存储器容量可以设置为0、8、16、32、64、100、132或164KB，对于计算能力8.6和8.9的设备，可以设置为零、8、16,32、64或100KB。 5) 应用程序可以使用cudaFuncSetAttribute（）设置划分，即首选共享内存容量。
cudaFuncSetAttribute（内核名称，cudaFuncAttributePreferredSharedMemoryCarveout，雕刻）；
API可以将分割指定为计算能力为8.0的设备支持的最大共享内存容量164 KB的整数百分比，计算能力为8.6和8.9的设备支持共享内存容量8.7和100 KB，也可以指定为以下值之一：｛cudaSharedmemCarveoutDefault、cudaShareDMemCarveutMaxL1或cudaSharedMCarveoutMaxShared。使用百分比时，将分割四舍五入到最近支持的共享内存容量。例如，对于计算能力为8.0%的设备，50%将映射到100 KB的分割，而不是82 KB的分割。考虑设置cudaFuncAttributePreferredSharedMemoryCarveout驱动程序的提示；如果需要，驱动程序可以选择不同的配置。 6) 具有计算能力8.0和8.7的设备允许单个线程块寻址高达163KB的共享存储器，而具有计算能力8.6和8.9的设备允许高达99KB的共享内存。依赖于每个块超过48KB的共享内存分配的内核是特定于体系结构的，并且必须使用动态共享内存，而不是静态大小的共享内存阵列。这些内核需要通过使用cudaFuncSetAttribute（）来设置cudaFuncAttributeMaxDynamicSharedMemorySize来显式选择加入；请参阅Volta体系结构的共享内存。 请注意，每个线程块的最大共享内存量小于每个SM的最大可用共享内存分区。线程块未可用的1 KB共享内存将保留给系统使用。 7) |
说点什么...