背景

在运行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的warp。

Case2

可以看到上图，kernel在时间线上GPU占用率分成了好几段，联系Case1的分析，当SM上有多个Block，SM优先运行同一个block的warp，所以一共应该分成了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，SM优先运行同一个block的warp。
3. 当kernel的block数<=SM数时，GPU会把各个block分配到独立的SM上，并同时运行。
4. 当kernel的block数>SM数时，GPU会把不同的block分配到同一个SM上，有多个block的SM，会优先运行同一个block的warp，可能出现阶梯状的占用率。

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

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

后记

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

参考

1. https://arxiv.org/pdf/2011.03897