（转载）GPU资源占用与利用率

来源：zhuanlan.zhihu.com/p/353410111

SM上的内存资源是有限的，如果每个线程占用的内存资源过多则一个SM上同时可执行的线程数就会减少。同理，如果每个线程块占用的内存资源过多，则一个SM上同时执行的线程块个数也会减少。

对于Tesla M6而言，主要内存资源的可用大小如下。
1） 寄存器：65536 per SM，65536 per Block
2） 共享内存：96KB per SM，48KB per Block
3） L2 Cache：2MB
线程会占用寄存器资源，占用过多会影响活跃的Warp数量；Block会占用共享内存资源，占用的过多会影响活跃的Block数量。对于Tesla M6而言，计算能力集5.2做出的限制如下。
1） 每个Warp有32个Thread
2） 每个Block最多可以包含1024个Thread
3） 每个SM最多可以并发64个Warp
4） 每个SM最多可以并发2048个Thread
5） 每个SM最多可以并发32个Block
6） 每个Thread最多使用255个寄存器
7） Warp分配粒度为4，寄存器以Warp为粒度分配，分配单元大小为256字节
8） 共享内存分配单元大小为256字节

核函数对SM利用率的计算公式为：

举两个极端例子。

1）每个Thread占用255个寄存器，每个Block内有256个线程，且Block不占用共享内存。此时一个Block占用65536个寄存器，由于SM内寄存器数量限制，则一个SM只能并发一个Block。一个Block有256个线程，转换为Warp为8个Warp。该情况下的SM Occupany = 8/64 = 12.5%，可见SM的利用率极低。
对于Tesla M6，为了将SM利用率达到100%，在块数够多的情况下，则：
a） 每个Block有1024个Thread时，每个Thread占用的寄存器数需小于等于32
b） 每个Block有512个Thread时，每个Thread占用的寄存器数需小于等于32
c） 每个Block有256个Thread时，每个Thread占用的寄存器数需小于等于32
d） 每个Block有128个Thread时，每个Thread占用的寄存器数需小于等于32
从以上信息可以看出，对于同一个GPU而言，无论Block中有多少个线程，为了最大化SM利用率，则每个Thread占用的最大寄存器数量是固定值。

如何查看核函数占用的寄存器数？

在编译选项中添加--ptxas-options=-v或者-Xptxas -v，则编译时会输出ptxas info信息，这些信息中包括核函数占用的寄存器数量、共享内存字节数以及常量内存字节数。
比如以下信息中，展示了kernelVelAddReal核函数占用了10个寄存器，348字节的常量内存，0字节的共享内存。

█ 注意，核函数的形参被存储在常量内存中，所以即使开发者没有主动使用常量内存，核函数也会占用一定数目的常量内存。（threadIdx.x这些变量存放在哪，寄存器里面吗？）
█ 注意，同一个核函数在Debug和Release模式下资源占用情况不一样，Release模式下编译器进行了优化。
█ 注意，一个空的核函数会占用2 registers，32bytes cmem[0]。（该信息不同平台不同版本可能不一样）

2）每个Thread占用0个寄存器，每个Block有个128线程，且每个Block占用48KB共享内存。由于SM内共享内存大小限制，此时一个SM只能并发2个Block，即可并发2*128/32=8个Warp，该情况下的SM Occupany = 8/64 = 12.5%，可见SM的利用率极低。
对于Tesla M6，为了将SM利用率达到100%，在块数够多的情况下，则：
a） 每个Block有1024个Thread时，每个Block占用共享内存需小于等于48KB
b） 每个Block有512个Thread时，每个Block占用共享内存需小于等于24KB
c） 每个Block有256个Thread时，每个Block占用共享内存需小于等于12KB
d） 每个Block有128个Thread时，每个Block占用共享内存需小于等于6KB
从以上信息可以看出，对于同一个GPU而言，Block内的Thread数越多，每个Block可占用的共享内存上限越大。

█ 注意，SM利用率最高并不代表性能最优，影响程序性能的因素有很多，SM利用率只是其中一个方面。