背景
在运行kernel时,需要设置kernel的Grid、Block值,我们一般都希望充分利用GPU的资源,达到最优性能。那么有如下问题:
- Grid、Block与SM占用率是什么关系?
- 如果想性能最优,该如何设置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%。分析:
- 在SM上的Block数= 93% x 108 x (64 / 32) = 200.88,忽略误差,与kernel的block数一致,都是200。
- 说明这一阶段kernel的所有block都已经分配到SM上。(对这里的计算逻辑和概念不够清楚的同学,如ComputeInFight等,可以翻看前文)。
继续观察kernel后半段的GPU占用率,发现ComputeInFight变成了43%,分析:
- kernel后半段时间在SM上的block数: 43% x 108 x (64 / 32) = 92,说明后半段在SM上的Block数是92个。
- 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占用率逐步下降。分析:
- 本例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个wave后,将按照每次递减108个block的逻辑进行,直到第10次减为0,这与图中现象一致,说明这里的分析逻辑是合理的。
其他Case
- 当单个block的线程数超过1024时,实际运行时会限制乘积为1024,这是GPU的限制。
- 当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时,仍然保持较高的占用率。
结论
- 单个block的线程数无法超过1024,即便代码配置超过,实际运行时还是会限制乘积为1024,这是GPU的限制。
- 一个SM上同一时刻只能运行一个block,GPU同时运行的block数不会超过SM总数。
- 当kernel的block数<=SM数时,GPU会把各个block分配到独立的SM上,并同时运行。
- 当kernel的block数>SM数时,GPU会把不同的block分配到同一个SM上,但不能同时运行。
a. kernel分成多个wave运行,wave数 = ceil(grids/SMs),SM占用率会逐步释放
所以,回到最开始的问题,如何设置Grid/Block值最合理?站在最大化GPU利用率的角度:
- Block thread数尽可能地接近1024,利用单个SM的计算能力。
- Grid值应尽可能整除SM数,避免Tail Effect。
后记
- 关于GPU利用率的资料很少,本文仅是从实验观测得出的结论,若有错误,请不吝指出。
参考
- https://arxiv.org/pdf/2011.03897