《CUDA并行程序设计-GPU编程指南》读书笔记--(1)线程网格、线程块以及线程

线程网格、线程块以及线程

SPMD模型

__global__ void addKernel(int *c, const int *a, const int *b)
{
    int i = threadIdx.x;
    c[i] = a[i] + b[i];
}

__global__前缀是告诉编译器在编译这个函数的时候生成的是GPU代码而不是CPU代码，并且这段GPU代码在CPU上是全局可见的。

我们可以将线程ID用作数组的下标对数组进行访问。线程0中threadIdx.x值为0，线程1的为1，依此类推，线程127中的threadIdx.x值为127。

CPU和GPU有各自独立的内存空间，因此在GPU代码中，不可以直接访问CPU端的参数，反过来在CPU代码中，也不可以直接访问GPU端的参数。

为了传递一个数据集到GPU端进行计算，我们需要使用cudaMalloc与cudaFree来申请 和释放显存，然后再使用cudaMemcpy将数据集从CPU端复制到GPU端，这样，才可以开始计算。

---

执行

线程都是以每32个一组，当所有32个线程都在等待诸如内存读取这样的操作时，它们就会被挂起。术语上，这些线程组叫做线程束(32个线程)或半个线程束(16个线程)。

将这128个线程分成4组，每组32个线程。首先让所有的线程提取线程标号，计算得到数组地址，然后发出一条内存获取的指令。接着下一条指令是做乘法，但这必须是在从内存读取数据之后。由于读取内存的时间很长，因此线程会挂起。当这组中的32个线程全部挂起，硬件就会切换到另一个线程束。

当线程束0由于内存读取操作而挂起时，线程束1就成为了正在执行的线程束。GPU一直以此种方式运行直到所有的线程束到成为挂起状态。

当连续的线程发出读取内存的指令时，读取操作会被合并或组合在一起执行。由于硬件在管理请求时会产生一定的开销，因此这样做将减少延迟(响应请求的时间)。由于合并，内存读取会返回整组线程所需要的数据，一般可以返回整个线程束所需要的数据。