CUDA进阶第七篇-如何获得核函数最佳线程块配置

最近在和几个小伙伴们一起翻译《cuda-c-programming-guide》，我选了chapter 5，收获良多。翻译到5.2.3.1 Occupancy Calculator的时候，正好跟现在手里的任务有些重合，就做了些实验，分享给大家。有说的不对的地方欢迎拍砖。

这一小节的内容很简单，如何获得核函数最佳线程块大小配置，可以用几个API，根据kernel内寄存器和共享内存的使用量，计算出block大小，使SM占用率最大。

内容很简单，译文还没校准完，大家可以先读英文。

我之前是不会这个的，所以我是用暴力循环的方式得出速度最快时的block大小，即以32为step，测block大小从32-1024时核函数的性能，但是这种方式比较麻烦，且平台和数据量都有一定的影响，不方便集成到系统中，所以只是实验的时候玩一玩。今天翻译到这个的时候，在想，这个occupancy方式是否真的好用，就做了个实验，并得出了一些经验行的结论。

代码里的核函数太简单了，我换了个复杂点的核函数，代码如下：

__global__ static void _bias_reluKer(float *data, float *bias, const int height, 
                                const int width)
{
    int bidx = blockIdx.x;
    int tidx = threadIdx.x;
    int span = blockDim.x;
    int i;

    if(bidx < height)
    {
        // copy data to shared memory
        i = tidx;
        for( ; i < width; i += span ) {
            float tmp = data[i + bidx * width] + bias[i]; // bias
            if(tmp < 0) tmp = 0.0f; 
            data[i + bidx * width] = tmp;
        }
    }
} 

cudaOccupancyMaxPotentialBlockSize()的结果是最佳blocksize是1024。

循环100遍，得到了总时间，实验结果如下：

arrayCount = 819200
OccupancyMax Out:blockSize=1024, gridSize = 800
32                                      : 8.022000(ms)
64                                      : 5.382000(ms)
704                                     : 5.348000(ms)
736                                     : 5.308000(ms)
768                                     : 5.249000(ms)
192                                     : 5.199000(ms)
800                                     : 5.194000(ms)
224                                     : 5.167000(ms)
832                                     : 5.145000(ms)
864                                     : 5.093000(ms)
896                                     : 5.050000(ms)
928                                     : 4.999000(ms)
352                                     : 4.964000(ms)
960                                     : 4.939000(ms)
544                                     : 4.917000(ms)
992                                     : 4.887000(ms)
416                                     : 4.880000(ms)
384                                     : 4.877000(ms)
576                                     : 4.873000(ms)
608                                     : 4.814000(ms)
448                                     : 4.808000(ms)
672                                     : 4.792000(ms)
640                                     : 4.776000(ms)
OccupancyMax                            : 4.774000(ms)
1024                                    : 4.763000(ms)
256                                     : 4.742000(ms)
480                                     : 4.738000(ms)
96                                      : 4.662000(ms)
512                                     : 4.612000(ms)
320                                     : 4.541000(ms)
128                                     : 4.507000(ms)
288                                     : 4.492000(ms)
160                                     : 4.435000(ms)

从实验结果我们可以得到结论：cudaOccupancyMaxPotentialBlockSize()函数确实是有效的，blocksize=1024和最佳blocksize=160的性能才差(4.77-1.43)/100 = 0.0034ms，可以直接忽略了。

但是，是否每个核函数都有必要这么做呢？因为这种方式并不流行，起码我CUDA编程5年以来第一次见，看caffe和kaldi的源码里好像也都没用这个函数。这么多年我都是直接用经验值，block=128 or 256。然后我又试几个核函数，结果发现……：

当核函数的计算不是十分复杂时，block=128 or 256 的性能跟Occupancy Calculator的结果差微乎其微，所以，还是直接用经验值吧……