背景
做GPU性能优化时，一种思路是先找到GPU占用率低的环节，然后做针对性优化，把GPU占用率提上去，那么问题就来了：
- 哪些指标可以表征GPU占用率？
- 如何测量GPU的占用率？
nsys profile gpu-metrics-devices
一番搜索后，发现nsys从 2021.2.4开始profile时开始支持gpu-metrics-devices参数，能够在profile时，按照一定频率记录GPU指标，参数如下：
实测有的版本参数是gpu-metrics-device，使用时可以都试下。在添加gpu-metrics-devices后，用nsys UI打开，会展示GPU的指标如图：
指标很多，包括SM使用率，nvlink带宽等等，文档都有介绍：
重点关注SM相关的指标：
- SM active：SM整体的使用率
- SM Warp Occupancy：SM warp占用情况
a. Compute Warps in Flight：SM中真正处于计算的warp占比
b. Unallocated Warps in Flight：活跃SM中，未使用的warp占比
如何理解SM这几个指标呢？谈谈个人的理解：
首先，对于特定GPU而言，SM数量、单个SM最大warp数都是固定的。比如A100一共有108个SM，单个SM最大warp数是64，换算成thread，就是2048个thread。
- SM active：按文档翻译是“SM至少有一个线程束（warp）在执行（分配在SM上）的周期数与总周期数的比率，以百分比表示。值为0表示所有SM都处于空闲状态（没有线程束在执行）。50%的值可能表示一个范围，从所有SM在采样期间50%的时间处于活跃状态，到50%的SM在整个采样期间100%活跃这两种极端情况之间的任何状态”。
a. 个人理解就是只要SM中有一个warp在运行，就会将该SM作为活跃SM，用活跃SM除以SM总数，得到SM active值，然后在周期内计算SM active的平均值。 -
Compute Warps in Flight：SM中真正在运行的warp的占比，假如GPU一共M个SM，每个SM支持N个warp，当前所有SM中真实运行的warp数为K，则该值为K /（M x N）。
-
Unallocated Warps in Flight：与Compute Warps in Flight相对应，含义是活跃的SM中未使用的warp的比率，计算方式为，设活跃的SM数为Q，则该值为（NxQ-K）/（M x N）。
理论上SM active = Compute Warps in Flight + Unallocated Warps in Flight，因为后两者的和就是活跃SM所有warp数除以warp总数，等于活跃SM占比。
测试
按照以上理解，我们写个程序测试下效果：
#include <stdio.h>
#include <iostream>
#include <chrono>
#include <vector>
#include <cuda_runtime_api.h>
#include <algorithm>
#include <numeric>
__global__ void kernel_test1(float *result)
{
unsigned int idx = threadIdx.x + blockDim.x * blockIdx.x;
for (int i = 99999; i >= 0; --i)
{
if (idx == 0)
{
result[0] = i;
}
}
}
void test(){
float *d_result;
cudaMalloc((void **)&d_result, 2 * sizeof(float));
// warm up
for (int i = 0; i < 10; ++i)
{
kernel_test1<<<dim3(1024, 1, 1), dim3(32, 1, 1)>>>(d_result);
}
// 按不同的thread数运行kernel
for (int j=32;j<1025;j*=2){
for (int i = 0; i < 1024; ++i)
{
kernel_test1<<<dim3(i, 1, 1), dim3(j, 1, 1),0>>>(d_result);
}
}
int main()
{
test();
cudaDeviceSynchronize();
return 0;
}
在A100上nsys profile结果为：
看整体趋势，在循环前期SM的占用率较低，后期占用变高。这是因为代码中的kernel使用的thread逐渐变多导致。观察其中一个kernel：
可以发现：
- 几乎每个block都会被分配到了一个独立的SM，因为SM active为43%，在周期内kernel平均占用了43% x 108=46.44个SM，与kernel grid几乎一致。
- 每个active SM中存在大量的未分配的warp，分配与未分配比例为5：37。
其实可以根据SM active、kernel参数、GPU硬件配置推测出Compute Warps in Flight与Unallocated Warps in Flight的值。以上为例子，推导过程如下：
- kernel所需总warp数：KernelWarps = grid x block / 32 = 46 x 256 / 32= 368
- ComputeWarpsInFlight = kernelWarps / GPUMaxWarps = 368 / (108 x 64) = 5.3%
- active SM支持的总warp数：ActiveSMWarps = SMs x SMActive x 64 = 108 x 43% x 64 = 2972
- UnallocatedWarpsInFlight = (ActiveSMWarps – KernelWarps）/ GPUMaxWarps = (2972 – 368)/(108 x 64)=37.6%
可以看到ComputeWarpsInFlight、UnallocatedWarpsInFlight与nsys的测量值基本一致，都是5:37，从侧面反应我们对以上指标的理解是合理的。
后记
- 本文内容仅代表个人理解，若有错误请不吝指出。
参考
- https://developer.nvidia.com/blog/measuring-the-gpu-occupancy-of-multi-stream-workloads/
- https://docs.nvidia.com/nsight-systems/UserGuide/index.html