GPU编程自学5 —— 线程协作

深度学习的兴起，使得多线程以及GPU编程逐渐成为算法工程师无法规避的问题。这里主要记录自己的GPU自学历程。

目录

• 《GPU编程自学1 —— 引言》
• 《GPU编程自学2 —— CUDA环境配置》
• 《GPU编程自学3 —— CUDA程序初探》
• 《GPU编程自学4 —— CUDA核函数运行参数》
• 《GPU编程自学5 —— 线程协作》
• 《GPU编程自学6 —— 函数与变量类型限定符》
• 《GPU编程自学7 —— 常量内存与事件》
• 《GPU编程自学8 —— 纹理内存》
• 《GPU编程自学9 —— 原子操作》
• 《GPU编程自学10 —— 流并行》

五、 线程协作

5.1 并行程序块的分解

首先回顾我们之前实现的矢量相加程序：

// 核函数：每个线程负责一个分量的加法
__global__ void addKernel(int *c, const int *a, const int *b)
{
    int i = threadIdx.x; // 获取线程ID
    c[i] = a[i] + b[i];
}

// 运行核函数，运行设置为1个block，每个block中size个线程
addKernel << <1, size >> >(dev_c, dev_a, dev_b);

通过前面小节，我们知道一个Block中的可开辟的线程数量是有限的(不超过1024)。

如果矢量特别长，上面的操作是会出现问题的。于是我们可以采用多个线程块(Block)来解决线程不足的问题。 假如我们设定每个线程块包含128个线程，则需要的线程块的数量为 size / 128。 为了避免不能整除带来的问题，我们可以稍微多开一点 (size + 127) / 128，但需要增加判断条件来避免越界。

// 核函数：每个线程负责一个分量的加法
__global__ void addKernel(int *c, const int *a,