如何理解GPU Kernel Grid/Block与SM占用率的关系?什么是Tail Effect?

背景

在运行kernel时,需要设置kernel的Grid、Block值,我们一般都希望充分利用GPU的资源,达到最优性能。那么有如下问题:

  1. Grid、Block与SM占用率是什么关系?
  2. 如果想性能最优,该如何设置Grid、Block值?

Grid、Block与SM占用率的关系

既然要探索二者的关系,直接写代码实验即可,有以下代码:

#include <iostream>
#include <cuda_runtime.h>

__global__ void myKernel(int *data) {
    const int bid = blockIdx.z * gridDim.x * gridDim.y + blockIdx.y * gridDim.x + blockIdx.x;
    const int tid = threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.x * blockDim.y;
    int temporary = data[0];
    for (int i = 0; i < 10000; ++i) {
        temporary *= temporary;
    }
    data[1] = temporary;
}

int main() {
    int *deviceData1;
    cudaMalloc((void**)&deviceData1, 2 * sizeof(int));
    cudaMemset(deviceData1, 0, 2 * sizeof(int));

    // warm up
    for(int i=0;i<200;i++){
        myKernel<<<dim3(32,1,1), dim3(1024,1,1)>>>(deviceData1);
    }
    cudaDeviceSynchronize();

    auto grids={
        dim3(1,1,1),
        dim3(10,1,1),
        dim3(100,1,1),
        dim3(107,1,1),
        dim3(108,1,1),
        dim3(109,1,1),
        dim3(200,1,1),
        dim3(216,1,1),
        dim3(1024,1,1),

        dim3(1,2,1),
        dim3(10,2,1),
        dim3(100,2,1),
        dim3(107,2,1),
        dim3(108,2,1),
        dim3(109,2,1),
        dim3(200,2,1),
        dim3(216,2,1),
        dim3(1024,2,1),

        dim3(1,1,2),
        dim3(10,1,2),
        dim3(100,1,2),
        dim3(107,1,2),
        dim3(108,1,2),
        dim3(109,1,2),
        dim3(200,1,2),
        dim3(216,1,2),
        dim3(1024,1,2),

        dim3(1,10,2),
        dim3(10,10,2),
        dim3(100,10,2),
        dim3(107,10,2),
        dim3(108,10,2),
        dim3(109,10,2),
        dim3(200,10,2),
        dim3(216,10,2),
        dim3(1024,10,2),
        };

    for (auto grid:grids){
        myKernel<<<grid, dim3(32,1,1)>>>(deviceData1);
        myKernel<<<grid, dim3(32,2,1)>>>(deviceData1);
        myKernel<<<grid, dim3(32,2,2)>>>(deviceData1);
        myKernel<<<grid, dim3(32,4,2)>>>(deviceData1);

        myKernel<<<grid, dim3(128,1,1)>>>(deviceData1);
        myKernel<<<grid, dim3(128,2,1)>>>(deviceData1);
        myKernel<<<grid, dim3(128,2,2)>>>(deviceData1);
        myKernel<<<grid, dim3(128,4,2)>>>(deviceData1);
    }

    cudaDeviceSynchronize();

    cudaFree(deviceData1);
    std::cout << "completed." << std::endl;
    return 0;
}

测试环境是A100,一共有108个SM,每个SM最多支持64个warp。perf后抽几个典型的case分析:

Case1


图中kernel每个block的warp数:128x4x2/32=32,观察到时间线上,kernel前半段的ComputeInFight是93%。分析:

  1. 在SM上的Block数= 93% x 108 x (64 / 32) = 200.88,忽略误差,与kernel的block数一致,都是200。
  2. 说明这一阶段kernel的所有block都已经分配到SM上。(对这里的计算逻辑和概念不够清楚的同学,如ComputeInFight等,可以翻看前文)。

继续观察kernel后半段的GPU占用率,发现ComputeInFight变成了43%,分析:

  1. kernel后半段时间在SM上的block数: 43% x 108 x (64 / 32) = 92,说明后半段在SM上的Block数是92个。
  2. kernel一共只有200个block,不可能先运行了200 block,又运行92 block。猜测前半段虽然block被分配到了SM上,但并未运行,所以先运行了108个,再运行剩余的92个block,这样就能完美解释了。

进一步可以猜测:单个SM在某一时刻只能运行一个block,即便SM上有多个Block。观察更多的kernel perf数据发现都服从这一规律。

Case2


可以看到上图,kernel在时间线上GPU占用率分成了好几段,联系Case1的分析,每时刻只能有SM数量个block在运行,所以一共应该分成了ceil(1024/108)=10段,即10个wave。这与上图是完全对应的。

再看kernel刚开始的时间段,SM ComputeInFlight是100%,而kernel一个block warp数是32x4x2/32=8,理论上单个block能在SM上的ComputeInFlight:8/64=12.5%。而图中的100%目前远比12.5%高,说明在开始阶段GPU会最大程度地把kernel预期的Block都分到了SM上。

观察后续过程,前2个wave ComputeInFlight一直是100%,后面GPU占用率逐步下降。分析:

  1. 本例GPU最多支持的kernel的block数:108×64/8=864,每个wave会消耗了108个block。在第一次wave后,还剩1000-108=892>864,所以第二个wave还是100%的 ComputeInFlight,第二次wave后,剩余1000-108×2=808<864,所以在第三个wave的时候,ComputeInFlight就会低于100%。
  2. 在第2个wave后,将按照每次递减108个block的逻辑进行,直到第10次减为0,这与图中现象一致,说明这里的分析逻辑是合理的。

其他Case

  1. 当单个block的线程数超过1024时,实际运行时会限制乘积为1024,这是GPU的限制。
  2. 当kernel的block数<=SM数时,GPU会把各个block分配到独立的SM上,并同时运行。

Tail Effect

观察以上case可以发现,在kernel末尾的GPU利用率很低,这是因为grid与GPU的SM数不匹配,导致最后的wave对SM的利用率不够,称为Tail Effect,详细可以查看论文:https://arxiv.org/pdf/2011.03897。


那么如何避免tail effect呢?

主要思路还是让grid尽可能整除SM数,这样便能在tail时,仍然保持较高的占用率。

结论

  1. 单个block的线程数无法超过1024,即便代码配置超过,实际运行时还是会限制乘积为1024,这是GPU的限制。
  2. 一个SM上同一时刻只能运行一个block,GPU同时运行的block数不会超过SM总数。
  3. 当kernel的block数<=SM数时,GPU会把各个block分配到独立的SM上,并同时运行。
  4. 当kernel的block数>SM数时,GPU会把不同的block分配到同一个SM上,但不能同时运行。
    a. kernel分成多个wave运行,wave数 = ceil(grids/SMs),SM占用率会逐步释放

所以,回到最开始的问题,如何设置Grid/Block值最合理?站在最大化GPU利用率的角度:

  1. Block thread数尽可能地接近1024,利用单个SM的计算能力。
  2. Grid值应尽可能整除SM数,避免Tail Effect。

后记

  1. 关于GPU利用率的资料很少,本文仅是从实验观测得出的结论,若有错误,请不吝指出。

参考

  1. https://arxiv.org/pdf/2011.03897
本文链接:https://rainlin.top/archives/305
转载请注明转载自:https://rainlin.top
暂无评论

发送评论 编辑评论


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