16.7. 计算能力 8.x
16.7.1. 架构
流处理器（SM）包括：
- 在计算能力为 8.0 的设备中有 64 个 FP32 核心用于单精度算术运算，而在计算能力为 8.6、8.7 和 8.9 的设备中有 128 个 FP32 核心，
- 在计算能力为 8.0 的设备中有 32 个 FP64 核心用于双精度算术运算，在计算能力为 8.6、8.7 和 8.9 的设备中有 2 个 FP64 核心，
- 64 个 INT32 核心用于整数运算，
- 4 个混合精度的第三代张量核心，支持半精度（fp16）、
__nv_bfloat16、tf32、子字节和双精度（fp64）矩阵运算，适用于计算能力为 8.0、8.6 和 8.7 的设备（详情见 Warp 矩阵函数）， - 4 个混合精度的第四代张量核心，支持
fp8、fp16、__nv_bfloat16、tf32、子字节和fp64，适用于计算能力为 8.9 的设备（详情见 Warp 矩阵函数）， - 16 个专用函数单元用于单精度浮点超越函数，
- 4 个 warp 调度器。
SM 静态地将其 warps 分配给各自的调度器。然后，在每个指令发出时间，每个调度器为其分配的、准备好执行的一个 warp 发出一个指令，如果有的话。
一个 SM 拥有：
- 一个只读的常量缓存，由所有功能单元共享，用于加速从设备内存中的常量内存空间的读取，
- 一个统一的数据缓存和共享内存，总大小为 192 KB，适用于计算能力为 8.0 和 8.7 的设备（是 Volta 的 128 KB 容量的 1.5 倍）和 128 KB，适用于计算能力为 8.6 和 8.9 的设备。
共享内存从统一数据缓存中划分出来，可以配置为各种大小（见 共享内存 部分）。剩余的数据缓存用作 L1 缓存，并且也被纹理单元使用，该单元实现了在 纹理和表面内存 中提到的各种寻址和数据过滤模式。
16.7.2. 全局内存
全局内存的行为与计算能力 5.x 的设备相同（见 全局内存）。
16.7.3. 共享内存
类似于 Volta 架构，保留给共享内存的统一数据缓存量可以根据每个内核进行配置。对于 NVIDIA Ampere GPU 架构，计算能力为 8.0 和 8.7 的设备的统一数据缓存大小为 192 KB，计算能力为 8.6 和 8.9 的设备为 128 KB。计算能力为 8.0 和 8.7 的设备的共享内存容量可以设置为 0、8、16、32、64、100、132 或 164 KB，计算能力为 8.6 和 8.9 的设备可以设置为 0、8、16、32、64 或 100 KB。
应用程序可以使用 cudaFuncSetAttribute() 设置 carveout，即首选的共享内存容量。
cudaFuncSetAttribute(kernel_name, cudaFuncAttributePreferredSharedMemoryCarveout, carveout);
API 可以将 carveout 指定为计算能力 8.0 和 8.7 设备支持的最大共享内存容量 164 KB 和计算能力 8.6 和 8.9 设备的 100 KB 的整数百分比，或者以下值之一：{cudaSharedmemCarveoutDefault、cudaSharedmemCarveoutMaxL1 或 cudaSharedmemCarveoutMaxShared。当使用百分比时，carveout 将四舍五入到最接近的支持的共享内存容量。例如，对于计算能力为 8.0 的设备，50% 将映射到 100 KB 的 carveout，而不是 82 KB。设置 cudaFuncAttributePreferredSharedMemoryCarveout 被驱动程序视为提示；如有需要，驱动程序可能选择不同的配置。
计算能力为 8.0 和 8.7 的设备允许单个线程块寻址高达 163 KB 的共享内存，而计算能力为 8.6 和 8.9 的设备允许高达 99 KB 的共享内存。依赖于每个块超过 48 KB 的共享内存分配的内核是特定于架构的，并且必须使用动态共享内存而不是静态大小的共享内存数组。这些内核需要通过使用 cudaFuncSetAttribute() 显式选择加入，以设置 cudaFuncAttributeMaxDynamicSharedMemorySize；见 Volta 架构的共享内存。
请注意，每个线程块的最大共享内存量小于每个 SM 可用的最大共享内存分区。不提供给线程块的 1 KB 共享内存保留供系统使用。