CUDA C/C++ 从入门到入土第二步——并行的思维

CUDA C/C++ 从入门到入土第一步——让你的CUDA跑起来
在第一步中配置好CUDA环境后，我们就具备了运行CUDA代码的现实条件了！但是，but, however, 很多时候，我们看着一个kernel函数，依然看不懂它的功能是什么，以及面对着cpu版本的函数，我们想要用CUDA核函数实现同样功能时，总感觉无从下手。这其实是因为我们太不习惯并行化了。希望大家看完本文后，能有一丢丢拜托串行思维的枷锁，对并行思维有更深一点的体会。

1.背景

有一个很常见的比喻，CPU是研究生，核心数量少（常听到的4核、8核），每一个核心都很强大，可以做很复杂的计算；
GPU是小学生，核心数量庞大（上千甚至更多），但每一个核心都比较菜，主打的是一个简单、门槛低的活，人多力量大。
根据CPU和GPU的特点，我们不难理解，对于CPU来说，核心强大的功能让我们这个boss基本只要发出逻辑正确的指令即可，然后剩下的事让CPU这个高材生自己处理。因此非常的省心，我们对CPU具体的操作可以当一个甩手掌柜。
但作为数量众多，小学生能力的GPU核心来说，他们需要被指定，哪些核心，要以哪些方式做哪些事。所以GPU编程总是会更累，需要考虑更全面的。

2. CUDA核函数体现出来的并行思维

对于如何更好阅读、以及写一个好的CUDA核函数，我认为背后的思想是：不要脱离群众！！要有群众的视角！！！
怎么理解呢？我们已经知道，我们作为leader，手下打工人都是一群小学生。对于你雇佣的这些小学生，他们每一个人都等着你派活干。
请记住，你的每一个指令，他们每个人都听到了并会执行。
例如：如果你想让某个学生吃饭，如果你下达的指令是“吃饭”，那么所有的小学生都会开始吃饭。

怎么做到不脱离群众呢？我的建议是换位思考：
把视角转换到核函数中一个普通而自信的线程！！

核函数中，请以blockIdx,threadIdx的线程为参考系。
你是一个线程，你看了看你boss给你发的工作牌，上面写着blockIdx和threadIdx
blockIdx.x = 1
blockIdx.y = 1
threadIdx.x = 2
threadIdx.y = 0
请注意，这些数字没有任何的特殊地方。你就是众多拥有blockIdx和threadIdx的线程中的一份子。完全可以假设你的blockIdx和threadIdx是另外的数值。

如果你入戏很深，你应该能想象到，在你的上下左右，都还有着其他的线程，他们有着自己的blockIdx和threadIdx，根据下图，你可以把你工作的block当成一个车间。整个任务由的车间数量为gridDim，每个车间有blockDim数量的线程。

我们常见到以下代码：

    int tx = blockIdx.x * blockDim.x +blockIdx.x;
    int ty = blockIdx.y * blockDim.y + threadIdx.y;

首先我们简单介绍一下，在CUDA中，我们默认左上角为原点，横轴朝右为x正方向，纵轴朝下为y正方向。
所以我们带入线程视角后，不难发现，此处的tx,ty代表了线程在gird当中的实际坐标。
也就是说，每一个线程都有自己的blockIdx和blockIdx，所以都有自己的坐标tx，ty。每一个线程有单独的（tx,ty），但是通过核函数中上述两行代码，在每一个thread中都计算得出了自己的坐标，并存入自己的寄存器中。

3.以矩阵乘法为例 3.1 cpu实现

void cpu_matrix_mult(int *h_a, int *h_b, int *h_result, int m, int n, int k) {
    for (int i = 0; i < m; ++i) 
    {
        for (int j = 0; j < k; ++j) 
        {
            int tmp = 0.0;
            for (int h = 0; h < n; ++h) 
            {
                tmp += h_a[i * n + h] * h_b[h * k + j];
            }
            h_result[i * k + j] = tmp;
        }
    }
}

3.2 CUDA核函数

我们在这一块实现一个矩阵乘法的核函数：a（m*n) * b(n*k) = c (m*k)
首先我们分析矩阵相乘的过程：c矩阵中的(i,j)为a矩阵中i行和b矩阵中j列逐元素成绩累加。

在CUDA变成中我们常常用的一种手段是：

屏蔽掉矩阵外的线程
将c矩阵的坐标和线程坐标对应起来。
背后的原理是，在C语言中，若有一个指针指向数组，例如：
type K[K1][K2];
type *p = (type *)K;
则K[m][n] = p[m * K2 + n];

我们重新看上述分析：“c(i,j)为a矩阵中i行和b矩阵中j列逐元素乘积累加。”
不难想到，可以用线程计算出的坐标对应c矩阵的坐标。即c(ty,tx)：
对每一个转化后坐标为tx,ty的线程，需要让a矩阵中ty行与b矩阵中ty列逐元素相乘累加。

 __global__ void mul( int *a, int *b, int *c, int m,int n,int k)
{
	//计算出每个线程的坐标值
    int tx = blockIdx.x * blockDim.x + threadIdx.x;
    int ty = blockIdx.y * blockDim.y + threadIdx.y;
    
    int sum = 0;
    if (tx < k && ty < m)  //屏蔽掉矩阵外的线程，线程坐标和矩阵坐标才可以对应起来
    {
        for(int i = 0; i < n; i++) 
        {
            sum += a[ty*n+i] * b[i*k+tx];  //逐元素相乘并求和
        }
        c[ty*k+tx] = sum;
    }
}

在写/读代码的过程中，如果蒙圈了，就好好想一想，我就是其中一个普通的任意位置的线程，blockIdx和 threadIdx是我的身份牌，此时代码中我做的是什么操作？或者我应当进行什么操作？

CUDA C/C++ 从入门到入土 第二步——并行的思维

C/C++/C#相关栏目本月热门文章

CUDA C/C++ 从入门到入土第二步——并行的思维