CUDA计算能力，线程调度及内存管理分析
CUDA计算能力，线程调度及内存管理分析
Compute Capability 6.x
1. 架构
SM 包括：
- 64（计算能力 6.0）或 128（6.1 和 6.2）CUDA 内核用于算术运算，
- 16（6.0）或32（6.1和6.2）用于单精度浮点超越函数的特殊功能单元，
- 2 （6.0） 或 4（6.1 和 6.2）warp调度程序。
当 SM 被赋予要执行的扭曲时，它首先将它们分发到其调度程序之间。然后，在每个指令发出时间，每个调度程序为其分配的 warps 之一发出一条指令，该指令已准备好执行（如果有）。
SM 具有：
- 由所有功能单元共享的只读常量缓存，并加快从驻留在设备内存中的常量内存空间的读取速度，
- 统一的 L1/纹理缓存，用于从大小为 24 KB（6.0 和 6.2）或 48 KB （6.1）的全局内存中读取，
- 大小为 64 KB（6.0 和 6.2）或 96 KB（6.1）的共享内存。
统一的L1/纹理缓存也由纹理单元使用，该纹理单元实现了纹理和表面内存中提到的各种寻址模式和数据过滤。
还有一个由所有 SMS 共享的 L2 缓存，用于缓存对本地或全局内存的访问，包括临时寄存器溢出。应用程序可以通过检查l2CacheSize设备属性来查询 L2 缓存大小（请参阅设备枚举）。
缓存行为（例如，读取是同时缓存在统一的 L1/纹理缓存和 L2 中还是仅在 L2 中）可以使用加载指令的修饰符在每次访问的基础上进行部分配置。
2. 全局内存
全局内存的行为方式与计算能力为 5.x 的设备相同（请参阅全局内存）。
3. 共享内存
共享内存的行为方式与计算能力为 5.x 的设备相同（请参阅共享内存）。
计算能力 7.x
1. 架构
SM 包括：
- 64 个 FP32 内核，用于单精度算术运算，
- 32 个 FP64 内核，用于双精度算术运算，34
- 64 个用于整数数学的 INT32 内核，
- 用于深度学习矩阵运算的 8 个混合精度张量核
- 16个用于单精度浮点超越函数的特殊功能单元，
- 4 个warp调度程序。
SM 在其调度程序之间静态分布其warp。然后，在每个指令发出时间，每个调度程序为其分配的 warps 之一发出一条指令，该指令已准备好执行（如果有）。
SM 具有：
- 由所有功能单元共享的只读常量缓存，并加快从驻留在设备内存中的常量内存空间的读取速度，
- 统一数据缓存和共享内存，总大小为 128 KB（伏特）或 96 KB（图灵）。
共享内存从统一数据缓存中分区出来，并且可以配置为各种大小（请参阅共享内存）。其余数据缓存用作 L1 缓存，也由实现纹理和表面内存中提到的各种寻址和数据过滤模式的纹理单元使用。
2. 独立的线程调度
Volta 架构在 warp 中的线程之间引入了独立的线程调度，从而实现了以前不可用的 warp 内同步模式，并简化了移植 CPU 代码时的代码更改。但是，如果开发人员对以前硬件体系结构的warp同步性做出假设，这可能会导致参与执行代码的线程集与预期大不相同。
以下是关注的代码模式以及针对 Volta 安全代码的建议纠正措施。
- 对于使用 warp 内部函数 （__shfl*, __any, __all, __ballot） 的应用程序，开发人员必须将其代码移植到带有*_sync后缀的新、安全、同步的对应项。新的经线内部函数采用线程掩码，明确定义哪些通道（经线）必须参与经线内部函数。有关详细信息，请参阅变形投票函数和变形随机函数。
由于 CUDA 9.0+ 提供了内部函数，因此（如有必要）可以使用以下预处理器宏有条件地执行代码：
#if defined(CUDART_VERSION) && CUDART_VERSION >= 9000
// *_sync intrinsic
#endif
这些内部函数适用于所有架构，而不仅仅是Volta或Turing，在大多数情况下，单个代码库足以满足所有架构的需求。但是请注意，对于 Pascal 和更早的架构，掩码中的所有线程都必须在收敛中执行相同的 warp 内在指令，并且掩码中所有值的并集必须等于 warp 的活动掩码。以下代码模式在 Volta 上有效，但在 Pascal 或更早的体系结构上无效。
if (tid % warpSize < 16) {
...
float swapped = __shfl_xor_sync(0xffffffff, val, 16);
...
} else {
...
float swapped = __shfl_xor_sync(0xffffffff, val, 16);
...
}
__ballot(1)的替代项是__activemask() 。请注意，即使在单个代码路径中，扭曲中的线程也可能发散。因此，__activemask()与__ballot(1)可能仅返回当前代码路径上的线程子集。以下无效代码示例在data[i]大于threshold 时将位i设置output为 1。__activemask() 用于尝试启用dataLen不是 32 的倍数的情况。
// Sets bit in output[] to 1 if the correspond element in data[i]
// is greater than 'threshold', using 32 threads in a warp.
for (int i = warpLane; i < dataLen; i += warpSize) {
unsigned active = __activemask();
unsigned bitPack = __ballot_sync(active, data[i] > threshold);
if (warpLane == 0) {
output[i / 32] = bitPack;
}
}
此代码无效，因为 CUDA 不保证warp仅在循环条件下发散。当由于其他原因发生背离时，扭曲中的不同线程子集将为相同的 32 位输出元素计算冲突的结果。正确的代码可能会使用__ballot_sync()非发散循环条件以及安全地枚举 warp 中参与阈值计算的线程集，如下所示。
for (int i = warpLane; i - warpLane < dataLen; i += warpSize) {
unsigned active = __ballot_sync(0xFFFFFFFF, i < dataLen);
if (i < dataLen) {
unsigned bitPack = __ballot_sync(active, data[i] > threshold);
if (warpLane == 0) {
output[i / 32] = bitPack;
}
}
}
发现模式演示了__activemask() 的有效使用情形。
- 如果应用程序具有扭曲同步代码，则需要在通过全局或共享内存在线程之间交换数据的任何步骤之间插入新的__syncwarp() warp 范围屏障同步指令。假设代码是同步执行的，或者从单独的线程读取/写入在扭曲中可见而不同步的假设是无效的。
- 尽管__syncthreads()一直记录为同步线程块中的所有线程，但 Pascal 和以前的架构只能在 warp 级别强制执行同步。在某些情况下，只要每个经线中至少有一些线程到达屏障，这就可以使屏障成功，而不会被每个线程执行。从Volta开始，CUDA内置__syncthreads()和PTX指令bar.sync（及其衍生产品）是按线程强制执行的，因此在块中的所有非退出线程到达之前不会成功。利用先前行为的代码可能会死锁，必须对其进行修改以确保所有未退出的线程都到达屏障。
2. __shared__ float s_buff[BLOCK_SIZE];
3. s_buff[tid] = val;
4. __syncthreads();
5.
6. // Inter-warp reduction
7. for (int i = BLOCK_SIZE / 2; i >= 32; i /= 2) {
8. if (tid < i) {
9. s_buff[tid] += s_buff[tid+i];
10. }
11. __syncthreads();
12. }
13.
14. // Intra-warp reduction
15. // Butterfly reduction simplifies syncwarp mask
16. if (tid < 32) {
17. float temp;
18. temp = s_buff[tid ^ 16]; __syncwarp();
19. s_buff[tid] += temp; __syncwarp();
20. temp = s_buff[tid ^ 8]; __syncwarp();
21. s_buff[tid] += temp; __syncwarp();
22. temp = s_buff[tid ^ 4]; __syncwarp();
23. s_buff[tid] += temp; __syncwarp();
24. temp = s_buff[tid ^ 2]; __syncwarp();
25. s_buff[tid] += temp; __syncwarp();
26. }
27.
28. if (tid == 0) {
29. *output = s_buff[0] + s_buff[1];
30. }
31. __syncthreads();
compute-saniter提供的racecheck 和 synccheck工具可以帮助查找违规行为。
为了在实施上述纠正措施时帮助迁移，开发人员可以选择加入不支持独立线程调度的 Pascal 调度模型。有关详细信息，请参阅应用程序兼容性。
3. 全局内存
全局内存的行为方式与计算能力为 5.x 的设备相同（请参阅全局内存）。
4. 共享内存
为共享内存保留的统一数据缓存量可基于每个内核进行配置。对于 Volta 架构（计算能力 7.0），统一数据缓存的大小为 128 KB，共享内存容量可以设置为 0、8、16、32、64 或 96 KB。对于图灵架构（计算能力 7.5），统一数据缓存的大小为 96 KB，共享内存容量可以设置为 32 KB 或 64 KB。与 Kepler 不同，驱动程序会自动为每个内核配置共享内存容量，以避免共享内存占用瓶颈，同时还允许在可能的情况下与已启动的内核并发执行。在大多数情况下，驱动程序的默认行为应提供最佳性能。
由于驱动程序并不总是知道完整的工作负载，因此应用程序有时提供有关所需共享内存配置的其他提示很有用。例如，共享内存使用很少或没有共享内存的内核可能会请求更大的剥离，以鼓励与需要更多共享内存的更高内核并发执行。新的cudaFuncSetAttribute() API 允许应用程序设置首选共享内存容量，或carveout作为支持的最大共享内存容量的百分比（Volta 为 96 KB，图灵为 64 KB）。
cudaFuncSetAttribute()与开普勒引入的传统cudaFuncSetCacheConfig() API 相比，放宽了首选共享容量的强制实施。旧版 API 将共享内存容量视为内核启动的硬要求。因此，将内核与不同的共享内存配置交错在一起，将不必要地序列化共享内存重新配置后的启动。使用新的 API，剥离被视为提示。如果需要执行函数或避免抖动，驱动程序可以选择不同的配置。
// Device code
__global__ void MyKernel(...)
{
__shared__ float buffer[BLOCK_DIM];
...
}
// Host code
int carveout = 50; // prefer shared memory capacity 50% of maximum
// Named Carveout Values:
// carveout = cudaSharedmemCarveoutDefault; // (-1)
// carveout = cudaSharedmemCarveoutMaxL1; // (0)
// carveout = cudaSharedmemCarveoutMaxShared; // (100)
cudaFuncSetAttribute(MyKernel, cudaFuncAttributePreferredSharedMemoryCarveout, carveout);
MyKernel <<<gridDim, BLOCK_DIM>>>(...);
除了整数百分比之外，还提供了几个方便的枚举，如上面的代码注释中列出。如果所选整数百分比未完全映射到支持的容量（SM 7.0 设备支持 0、8、16、32、64 或 96 KB 的共享容量），则使用下一个更大的容量。例如，在上面的示例中，最大 50 KB 的 96% 为 48 KB，这不是受支持的共享内存容量。因此，首选项向上舍入为 64 KB。
计算能力 7.x 设备允许单个线程块处理共享内存的全部容量：Volta 上为 96 KB，图灵上为 64 KB。依赖于每个块超过 48 KB 的共享内存分配的内核是特定于体系结构的，因此它们必须使用动态共享内存（而不是静态大小的数组），并且需要使用cudaFuncSetAttribute()如下方式显式选择加入。
// Device code
__global__ void MyKernel(...)
{
extern __shared__ float buffer[];
...
}
// Host code
int maxbytes = 98304; // 96 KB
cudaFuncSetAttribute(MyKernel, cudaFuncAttributeMaxDynamicSharedMemorySize, maxbytes);
MyKernel <<<gridDim, blockDim, maxbytes>>>(...);
否则，共享内存的行为方式与计算能力为 5.x 的设备相同（请参阅共享内存）。
