grid和block参数设置和地址计算注意事项

Posted on Views:

cuda中，核函数的参数block和grid的设置及其重要，符合硬件特点的设置可以提高硬件资源的利用效率。核函数编写中，设计到通用矩阵乘（GEMM）优化算法，grid和block组合下thread的排列和地址计算也容易让人绕的头晕。而且矩阵常见的地址计算方式和多维数组寻址还不一样；
原文主要参考，加上一些自己的理解。

block

同一个 block 中，连续的 32 个线程组成一个 warp，这 32 个线程每次执行同一条指令，也就是所谓的 SIMT，即使最后一个 warp 中有效的线程数量不足 32，也要使用相同的硬件资源，所以 block_size 最好是 32 的整数倍。

SM

与 block 对应的硬件级别为 SM，SM 为同一个 block 中的线程提供通信和同步等所需的硬件资源，跨 SM 不支持对应的通信，所以一个 block 中的所有线程都是执行在同一个 SM 上的，而且因为线程之间可能同步，所以一旦 block 开始在 SM 上执行，block 中的所有线程同时在同一个 SM 中执行（并发，不是并行），也就是说 block 调度到 SM 的过程是原子的。SM 允许多于一个 block 在其上并发执行，如果一个 SM 空闲的资源满足一个 block 的执行，那么这个 block 就可以被立即调度到该 SM 上执行，具体的硬件资源一般包括寄存器、shared memory、以及各种调度相关的资源。
这里的调度相关的资源一般会表现为两个具体的限制，Maximum number of resident blocks per SM 和 Maximum number of resident threads per SM，也就是 SM 上最大同时执行的 block 数量和线程数量。

Occupancy

提高SM的利用率，其中一个最简单的方法是让尽量多的线程同时在 SM 上执行，SM 上并发执行的线程数和SM上最大支持的线程数的比值，被称为 Occupancy，更高的 Occupancy 代表潜在更高的性能。对于挖掘一块gpu整体的性能（可能多个核函数同时跑），显然，一个 kernel 的 block_size 应大于 SM 上最大线程数和最大 block 数量的比值，否则就无法达到 100% 的 Occupancy。对于单个任务来说，不考虑整个gpu的挖掘，则不一定，具体可以自己尝试不同的block_size。
对应不同的架构，这个比值不相同，对于 V100 、 A100、 GTX 1080 Ti 是 2048 / 32 = 64，对于 RTX 3090 是 1536 / 16 = 96，所以为了适配主流架构，如果静态设置 block_size 不应小于 96。考虑到 block 调度的原子性，那么 block_size 应为 SM 最大线程数的约数，否则也无法达到 100% 的 Occupancy，主流架构的 GPU 的 SM 最大线程数的公约是 512，96 以上的约数还包括 128 和 256，也就是到目前为止，block_size 的可选值仅剩下 128 / 256 / 512 三个值。

还是因为 block 调度到 SM 是原子性的，所以 SM 必须满足至少一个 block 运行所需的资源，资源包括 shared memory 和寄存器，shared memory 一般都是开发者显式控制的，而如果 block 中线程的数量 * 每个线程所需的寄存器数量大于 SM 支持的每 block 寄存器最大数量，kernel 就会启动失败。

目前主流架构上，SM 支持的每 block 寄存器最大数量为 32K 或 64K 个 32bit 寄存器，每个线程最大可使用 255 个 32bit 寄存器，编译器也不会为线程分配更多的寄存器，所以从寄存器的角度来说，每个 SM 至少可以支持 128 或者 256 个线程，block_size 为 128 可以杜绝因寄存器数量导致的启动失败，但是很少的 kernel 可以用到这么多的寄存器，同时 SM 上只同时执行 128 或者 256 个线程，也可能会有潜在的性能问题。但把 block_size 设置为 128，相对于 256 和 512 也没有什么损失，128 作为 block_size 的一个通用值是非常合适的。

grid

确定了 block_size 之后便可以进一步确定 grid_size，也就是确定总的线程数量，对于一般的 elementwise kernel 来说，总的线程数量应不大于总的 element 数量，也就是一个线程至少处理一个 element，同时 grid_size 也有上限，为 Maximum x-dimension of a grid of thread blocks，目前在主流架构上都是 2^31 - 1，对于很多情况都是足够大的值。
GPU 一次可以调度 SM 数量 * 每个 SM 最大 block 数个 block，因为每个 block 的计算量相等，所以所有 SM 应几乎同时完成这些 block 的计算，然后处理下一批，这其中的每一批被称之为一个 wave。想象如果 grid_size 恰好比一个 wave 多出一个 block，因为 stream 上的下个 kernel 要等这个 kernel 完全执行完成后才能开始执行，所以第一个 wave 完成后，GPU 上将只有一个 block 在执行，GPU 的实际利用率会很低，这种情况被称之为 tail effect。

我们应尽量避免这种情况，将 grid_size 设置为精确的一个 wave 可能也无法避免 tail effect，因为 GPU 可能不是被当前 stream 独占的，常见的如 NCCL 执行时会占用一些 SM。所以无特殊情况，可以将 grid_size 设置为数量足够多的整数个 wave，往往会取得比较理想的结果，如果数量足够多，不是整数个 wave 往往影响也不大。

当然，这里还要注意查看设备SM的数量，设置为其的整数倍，能保证SM的负载比较均衡；

grid和block设置总结

综上所述，普通的 elementwise kernel 或者近似的情形中，block_size 设置为 128，grid_size 设置为可以满足足够多的 wave 就可以得到一个比较好的结果了。但更复杂的情况还要具体问题具体分析，比如，如果因为 shared_memory 的限制导致一个 SM 只能同时执行很少的 block，那么增加 block_size 有机会提高性能，因为共享内存和寄存器公用一块物理内存，共享内存用多了，给sm给线程分配的寄存器空间小了，线程数就会受到限制；如果 kernel 中有线程间同步，那么过大的 block_size 会导致实际的 SM 利用率降低，因为物理上不同线程束实际上是并发执行的，所以不同线程束之间势必存在等待的情况。
还可以参考这一篇也很有启发性：https://blog.csdn.net/quicmous/article/details/115250605

grid和block常见地址计算

多维数组寻址中，如果是二维的矩阵，可以看成是按行（内层）一个个地址读下去，让这个地址对应的线程干一件事儿。但是往往GEMM中，比如简单的矩阵相乘，一个线程里可能要操作不同两个矩阵中不同位置的元素相乘，那么，这里对grid和block的理解，或者按照什么方式去使用就很重要了。

行列方式计算索引

1
2
3
4
5
6
7
8
9
10
11
12
13
14
__global__ void matMul(float *d_A, float *d_B, float *d_C, int M, int N, int P) {
        int row = blockIdx.y * blockDim.y + threadIdx.y; //行
        int col = blockIdx.x * blockDim.x + threadIdx.x; //列

        if(row < M && col < P) {
                float sum = 0.0f;

                // compute the dot product for each row of A and col of B
                for(int i = 0; i < N; ++i) {
                        sum += d_A[row * N + i] * d_B[i * P + col];
                }
                d_C[row * P + col] = sum;
        }
}

上诉就是按照行列去计算索引，然后通过传入矩阵的维度，M N P方便遍历和计算, 达到按行和列相乘的目的，相当于通过行列索引分别的计算，使得当前对应的线程可以去索引不同矩阵ABC去做计算。当然，这里矩阵的维度M和N可以考虑用grid的维度替代。

多维数组方式计算计算索引

参考多维数组寻址即可，这里要说明的是，采取这种方式，相当于一个线程只拿当前计算的唯一地址索引。
此外，二维的排列还可以参考这种计算索引方式：
cuda_thread_index
实际上idx=iy*nx+ix代入拆开后和多维数组方式计算的规律差不多。

正在加载今日诗词....