如何优雅地测量GPU占用率?(nsys gpu-metrics-devices使用)

背景

做GPU性能优化时,一种思路是先找到GPU占用率低的环节,然后做针对性优化,把GPU占用率提上去,那么问题就来了:

  1. 哪些指标可以表征GPU占用率?
  2. 如何测量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相关的指标:

  1. SM active:SM整体的使用率
  2. 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。

  1. SM active:按文档翻译是“SM至少有一个线程束(warp)在执行(分配在SM上)的周期数与总周期数的比率,以百分比表示。值为0表示所有SM都处于空闲状态(没有线程束在执行)。50%的值可能表示一个范围,从所有SM在采样期间50%的时间处于活跃状态,到50%的SM在整个采样期间100%活跃这两种极端情况之间的任何状态”。
    a. 个人理解就是只要SM中有一个warp在运行,就会将该SM作为活跃SM,用活跃SM除以SM总数,得到SM active值,然后在周期内计算SM active的平均值。

  2. Compute Warps in Flight:SM中真正在运行的warp的占比,假如GPU一共M个SM,每个SM支持N个warp,当前所有SM中真实运行的warp数为K,则该值为K /(M x N)。

  3. 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:



可以发现:

  1. 几乎每个block都会被分配到了一个独立的SM,因为SM active为43%,在周期内kernel平均占用了43% x 108=46.44个SM,与kernel grid几乎一致。
  2. 每个active SM中存在大量的未分配的warp,分配与未分配比例为5:37。

其实可以根据SM active、kernel参数、GPU硬件配置推测出Compute Warps in Flight与Unallocated Warps in Flight的值。以上为例子,推导过程如下:

  1. kernel所需总warp数:KernelWarps = grid x block / 32 = 46 x 256 / 32= 368
  2. ComputeWarpsInFlight = kernelWarps / GPUMaxWarps = 368 / (108 x 64) = 5.3%
  3. active SM支持的总warp数:ActiveSMWarps = SMs x SMActive x 64 = 108 x 43% x 64 = 2972
  4. UnallocatedWarpsInFlight = (ActiveSMWarps – KernelWarps)/ GPUMaxWarps = (2972 – 368)/(108 x 64)=37.6%

可以看到ComputeWarpsInFlight、UnallocatedWarpsInFlight与nsys的测量值基本一致,都是5:37,从侧面反应我们对以上指标的理解是合理的。

后记

  1. 本文内容仅代表个人理解,若有错误请不吝指出。

参考

  1. https://developer.nvidia.com/blog/measuring-the-gpu-occupancy-of-multi-stream-workloads/
  2. https://docs.nvidia.com/nsight-systems/UserGuide/index.html
本文链接:https://rainlin.top/archives/285
转载请注明转载自:https://rainlin.top
暂无评论

发送评论 编辑评论


				
|´・ω・)ノ
ヾ(≧∇≦*)ゝ
(☆ω☆)
(╯‵□′)╯︵┴─┴
 ̄﹃ ̄
(/ω\)
∠( ᐛ 」∠)_
(๑•̀ㅁ•́ฅ)
→_→
୧(๑•̀⌄•́๑)૭
٩(ˊᗜˋ*)و
(ノ°ο°)ノ
(´இ皿இ`)
⌇●﹏●⌇
(ฅ´ω`ฅ)
(╯°A°)╯︵○○○
φ( ̄∇ ̄o)
ヾ(´・ ・`。)ノ"
( ง ᵒ̌皿ᵒ̌)ง⁼³₌₃
(ó﹏ò。)
Σ(っ °Д °;)っ
( ,,´・ω・)ノ"(´っω・`。)
╮(╯▽╰)╭
o(*////▽////*)q
>﹏<
( ๑´•ω•) "(ㆆᴗㆆ)
😂
😀
😅
😊
🙂
🙃
😌
😍
😘
😜
😝
😏
😒
🙄
😳
😡
😔
😫
😱
😭
💩
👻
🙌
🖕
👍
👫
👬
👭
🌚
🌝
🙈
💊
😶
🙏
🍦
🍉
😣
Source: github.com/k4yt3x/flowerhd
颜文字
Emoji
小恐龙
花!
上一篇
下一篇