CUDA编程中gridsize,blocksize大小怎么设置

基本概念

在 CUDA 中，GPU 的计算任务被组织成 Grid（网格） 和 Block（线程块）：

比如我们经常看到CUDA内核函数的调用方式是：

kernel<<<gridSize, blockSize>>>(...);

blockSize：每个 Block 包含的线程数，这是一个dim3的参数（blockDim.x, blockDim.y, blockDim.z）。

gridSize：整个 Grid 包含的 Block 数，这也是一个dim3的参数（gridDim.x, gridDim.y, gridDim.z）。

blockSize和gridSize都是一个dim3类型的参数，一般来说y，z都为1简化索引的复杂度，
例如：
blockSize = 256（1D）
blockSize = dim3(16, 16)（2D，共 256 线程）
blockSize = dim3(8, 8, 4)（3D，共 256 线程）
这几种表达方式都是合法的。

问题：这里面gridSize和blockSize具体应该怎么设置呢

gridSize和blockSize大小的设置

blcokSize大小的设置

1. Block所包含的线程数在软件上有最大值1024，即blockSize <1024。
2. Block线程之间的信息可以通信，而不同SM之间线程无法通信，所以SM的范围是比block的范围大的。即Block是在SM上的一个软件上抽象出来的概念,其所有的参数会受到SM硬件资源的限制。

blockSize大小设置范围

由1可以知道，如果不考虑任何优化手段，只希望代码能正常运行，blockSize可以设置1-1024的任何值。这是一个大前提。
但是blockSize大小的设置的好不好，对核函数的影响性能非常大。所以我们可以接着看看有哪些可以优化的手段。

blockSize大小设置优化方向

前面提到了block的概念是＜SM，所以会受到SM硬件的限制，具体的限制主要有：

1. SM最大寄存器的限制
2. SM最大shared memory限制
3. SM最大Block的限制
4. SM最大Thread的限制
5. 每个warp执行相同的指令

一、寄存器的限制

每个SM最多有64K个32bit的寄存器，每个线程最多255个寄存器，假设计算非常复杂需要用满寄存器的数量.那么这个SM最大运行的线程数量Thread = 65,536 / 256 = 256个。也就是说如果blockSize设置＞256时，有可能会出现寄存器不够的情况。

这种情况比较少见，所以这个优化手段可以了解一下，并不一定需要在实际代码中用到

二、shared memory的限制
shared memory只有64K，需要看看每个block的shared memory够不够用

三、SM的最大Block和最大Thread的限制
1.不同显卡的最大Block和最大Thread是有限制的，如3090，SM的最大block数量是16，最大的Thread数量是1536，为了保证线程能吃满的话，必须要保证每个block的线程数：
$blockSize＞1536/16=96$
, 因为最大的block数量是16，如果小于这个值，SM肯定有部分线程没有用到。
2.一个SM可以放多个Block，所以可以让blockSize设置为最大Thread的约数。这样可以平分吃掉所有的Thread。

四、每个warp执行相同的指令
每个warp有32个线程，为了保证同一个warp执行相同的指令，blockSize的数量最好设置为32的倍数。

总结

所以blockSize的设置最好为128,256。如果计算不太复杂（一般计算都不太复杂），可以设置为512。

gridSize

gridSize设置数量范围是 $[1,2^{31}-1]$,保证一次运算能吃满所有的SM。

参考资料:如何设置CUDA Kernel中的grid_size和block_size？