本系列相关文章：
为了能够更深入理解GPU Profiling，前面用了一整篇文章介绍了CUDA编程和GPU架构相关知识。如果对CDUA编程和GPU架构不太了解，建议先看看之前的文章。本篇文章将继续介绍GPU Profiling。
GPU利用率的局限性
大多数时候，我们都是通过nvidia-smi查询GPU利用率来了解GPU使用情况，如果该值比较高，我们就认为GPU使用得比较好。
GPU利用率的意义可以描述为：在一个周期时间内（1s或1/6s，依GPU产品而定），一个或多个核函数处于active的时间（eg：running）。 但是GPU利用率目前存在如下的局限性：
-
仅能够展示有核函数在用GPU资源，但用的怎么样无法知道。
-
无法告知用户有多少个SM在使用。
-
用户写的代码有多“忙”，到底是真忙还是假忙？
-
用户写的代码到底在干什么？比如当前是在做单精度浮点数计算（FP32）还是做双精度浮点数计算（FP64）还是在从内存中拷贝数据？
为了演示GPU利用率的局限性，接下来我们做一个实验，以下是一段示例CUDA代码Test_GPU_Utilization.cu：
$ cat Test_GPU_Utilization.cu
#include
#include
#include
const long long tdelay=100000000000000LL;
// 核函数定义，这个核函数所干的事有点类似于使这个GPU线程一直sleep
__global__ void dkern(){
long long start = clock64();
long long targetClock = start+tdelay;
while(clock64() < targetClock);
}
int main(int argc, char *argv[]){
dkern<<<1,1>>>(); // 调用核函数启动1个线程块，这个线程块包含1个线程
cudaDeviceSynchronize();
return 0;
}然后使用nvcc编译这段代码并执行：
# nvcc -o test-gpu-util Test_GPU_Utilization.cu
# ./test-gpu-util在另外的终端上使用nvidia-smi查看GPU利用率，可以看到GPU利用率为100%：
虽然GPU利用率显示了100%，但是真的就说明代码高效利用了这块GPU吗？我们来分析一下，我是在V100上运行这个程序的，前面的文章介绍过，V100使用的计算能力（Compute Capability,简称CC）为7.0，在CC 7.0中定义了每个SM最大能容纳的线程数为2048，且V100有84个SM，那么整张V100在极限情况下，一次能够容纳的线程数为84 * 2048 = 172032，而示例程序中，核函数只启动了1个线程块，这个线程块包含1个线程。可以看到，V100一次本来能够容纳172032个线程，现在却启动了1个线程，利用率也是100%，误差是不是很大？
DCGM Profiling指标
NVIDIA认识到单就GPU利用率这一个指标并不能完全说明GPU的使用情况，既然一个指标不能说明，那就多弄几个指标，从各个维度进行说明，DCGM Profiling相关指标就由此产生了。
下表是对DCGM提供的Profiling指标做一些解释，更多指标解释请参考文档：容器服务GPU监控2.0指标说明。
指标名称 |
解释 |
DCGM_FI_PROF_GR_ENGINE_ACTIVE |
|
DCGM_FI_PROF_SM_ACTIVE |
|
DCGM_FI_PROF_SM_OCCUPANCY |
|
DCGM_FI_PROF_PIPE_TENSOR_ACTIVE |
|
DCGM_FI_PROF_PIPE_FP64_ACTIVE |
|
DCGM_FI_PROF_PIPE_FP32_ACTIVE |
|
DCGM_FI_PROF_PIPE_FP16_ACTIVE |
|
DCGM_FI_PROF_DRAM_ACTIVE |
|
|
|
|
|
使用GPU监控2.0做Profiling
提交示例任务
现在我们仍然使用前面“GPU利用率的局限性”小节中提到的那段测试GPU利用率的代码，并制作成docker镜像，然后使用kubectl apply提交一个任务，任务的yaml如下：
apiVersion: batch/v1
kind: Job
metadata:
name: test-gpu-util
spec:
parallelism: 1
template:
metadata:
labels:
app: test-gpu-util
spec:
containers:
- name: test-gpu-util
image: registry.cn-beijing.aliyuncs.com/ai-samples/gpu-util-test:v0.1.1
resources:
limits:
nvidia.com/gpu: 1
restartPolicy: Never任务提交后，等待任务的pod处于Running：
# kubectl get po -owide
NAME READY STATUS RESTARTS AGE IP NODE NOMINATED NODE READINESS GATES
test-gpu-util-wx22z 1/1 Running 0 13m 10.2.1.131 cn-beijing.192.168.10.171 可以发现pod运行在节点cn-beijing.192.168.10.171上。
登录GPU监控2.0大盘
登录到GPU监控2.0大盘，具体操作如下：
-
在 集群列表 页面中，单击目标集群名称或者目标集群右侧 操作 列下的 详情 。
-
在集群管理页左侧导航栏中，选择 运维管理 > Prometheus监控 。
-
在 Prometheus监控 大盘列表页面，单击 GPU监控 页签，您分别可以看到 集群维度的GPU监控大盘 和 节点维度的GPU监控大盘 ， 点击“节点维度GPU监控大盘” 。
进入大盘后，通过左上角“GPUNode”选项框可以选择节点，因为pod运行在cn-beijing.192.168.10.171，那么选择节点cn-beijing.192.168.10.171。从大盘中可以看到该节点只有1张GPU卡，且已被分配。
分析大盘数据
为了更好的理解GPU Profiling效果，先做如下的说明：
-
当前节点使用的卡型为Tesla V100
-
从V100的白皮书可以知道Tesla V100拥有84个SM，能够使用的SM数为80
-
V100使用的计算能力是7.0版本，7.0版本定义的物理限制如下
在“Utilization”这一栏，GPU利用率（GPU Utilization）为100%（也就是执行nvidia-smi显示的值），内存带宽利用率（Memory Copy Utilization）为0，说明节点上的0号GPU卡没有发生与内存有关的操作，这与事实比较相符，因为test-gpu-util任务的核函数没有从内存读取数据或写入数据到内存。
在“Profiling”这一栏，SM Active表示当前GPU卡上有多少个SM处于“Active”状态，“Active”状态的意思只要这个SM有线程，那么就是"Active"的。任务test-gpu-util的核函数会一直运行，那么处于Active的SM数为80（一张GPU卡CUDA可用SM数） * 1.25% = 1，说明只有1个SM处于Active，符合预期，因为任务test-gpu-util的核函数只启动了一个线程块，且该线程块只有1个线程，只能在一个SM上运行。虽然test-gpu-util任务的GPU利用率达到100%，但在SM Active这个指标面前它无所遁形，80个SM只用了1个，说明GPU利用率是虚高。
SM Occupancy显示的是每张GPU卡的CUDA占用率，大盘显示0号GPU卡的占用率为0.0195%，即采样周期内每个SM的平均占用率为0.0195%，那么80个SM总的占用率为：80 * 0.0195% = 1.56%，由SM Active知道只有1个SM是被使用的，那么1.68%是属于这个SM的占用率，V100采用计算能力7.0版本，而计算能力7.0版本中规定每个SM最大可驻留的线程数为2048，那么这个SM当前驻留的线程数为2048 * 1.56% = 31.95，约为32个线程，符合预期。这里你可能有疑惑：任务test-gpu-util的核函数只启动了一个线程块且该线程块只有1个线程，为什么通过占用率计算出SM当前驻留的线程个数为32，而不是1呢？前面的文章中我们提到过：GPU的SM上是以线程束为单位执行的，而每个线程束包含32个线程，也就是说，SM要么不执行，如果要执行，至少执行32个线程。
从占用率这个指标也可以看出：test-gpu-util这个任务并没有高效利用GPU，因为总共可以驻留2048个线程，当前只驻留了32个线程。
由于test-gpu-util这个任务并没有执行与FP16、FP32、FP64等有关的运算，所以这些指标的面板显示数据均为0，从这些指标中可以了解到一个任务是“真忙”还是“假忙”：它到底是忙于做FP32类的计算？还是忙于做FP64类的计算？还是像test-gpu-util这个任务一样，做sleep样的动作？
同样，通过DRAM Active可以知道核函数是否存在对GPU内存操作的动作，由于DRAM Active显示为0，说明test-gpu-util这个任务并没有操作GPU内存。
结论
从监控大盘中可以得到对test-gpu-util这个任务的如下描述：
-
GPU利用率为100%，内存带宽利用率为0
-
80个SM只用了其中1个，其余79个SM处于空闲状态
-
SM能够容纳2048个线程，当前却只容纳了32个线程
-
FP32、FP64、FP16等计算基本没有
-
也没有对内存有存取操作
所以得出的结论是：test-gpu-util只是一个能让GPU利用率虚高的任务，没有任何实际意义。
总结
本篇文章只是通过一个很简单的核函数介绍了如何通过阿里云容器服务GPU监控2.0做GPU Profiling，但在实际的场景中会比这更为复杂，因为一个CUDA程序有几十或者上百个核函数（而test-gpu-util这个示例任务只有1个核函数），有时还会存在多个CUDA程序运行在同一张GPU卡上。但无论如何，这些指标对于了解GPU的使用状况是有帮助的。
