线程网格、线程块以及线程
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一直以此种方式运行直到所有的线程束到成为挂起状态。
当连续的线程发出读取内存的指令时,读取操作会被合并或组合在一起执行。由于硬件在管理请求时会产生一定的开销,因此这样做将减少延迟(响应请求的时间)。由于合并,内存读取会返回整组线程所需要的数据,一般可以返回整个线程束所需要的数据。