CUDA学习笔记（三）-线程通信

在上一章中，我们还看到了maxThreadPerBlock属性，它限制了每个块可以启动的线程数量。对于最新的gpu，它的值是512或1,024。同样，在第二种方法中，可以并行启动的最大块数被限制为65535
理想情况下，我们不是针对单个块启动多个线程，或者使用单个线程启动多个块，而是启动多个块，每个块并行地拥有多个线程(可以等于maxThreadPerBlock)。那么，假设你想要并行地启动N = 50,000个线程，在vector add的例子中，我们在上一章看到过。

gpuAdd<< <((N +511)/512),512 > >>(d_a,d_b,d_c)

每个块的最大线程是512个，因此块的总数是通过线程总数(N)除以512来计算的。但是如果N不是512的确切倍数，那么N除以512可能给出错误的块数，它比实际的计数少1。所以，为了得到块数的下一个最大的整数值，所以N应该加上511，然后除以512。
但是传统的写法当N大于65535时就会失效，所以我们应该怎么做呢？
还是以一段修改后的代码来进行讲解

#include "stdio.h"
#include<iostream>
#include <cuda.h>
#include <cuda_runtime.h>
//Defining number of elements in array
#define N 50000
__global__ void gpuAdd(int *d_a, int *d_b, int *d_c)
{
//Getting index of current kernel
int tid = threadIdx.x + blockIdx.x * blockDim.x;
while (tid < N)
{
d_c[tid] = d_a[tid] + d_b[tid];
tid += blockDim.x * gridDim.x;
}
}

这个内核代码类似于我们在上一章中所写的。它有两个修改。一个修改是在线程ID的计算中，第二个修改是在内核函数中包含while循环。线程ID计算的变化是由于并行地启动多个线程和块。可以将块和线程看作是一个二维矩阵，块数等于行数，列数等于每个块的线程数。我们将以三个块和三个线程/块为例，如下表所示：

如果我们想在所有线程中为这个绿色的线程创建一个惟一的索引就需要进行一个数学计算：`tid = threadIdx.x + blockIdx.x * blockDim.x;
例如，绿色的，threadIdx.x = 1, blockIdx.x = 2 , blockDim.x = 3 则 tid = 7.学习这个计算非常重要，因为它将在我们以后的代码中广泛使用。
代码中包含while循环是因为当N非常大时，由于前面描述的限制，线程总数不能等于N。因此，一个线程必须执行多个操作，这些操作由启动的线程总数分隔。这个值可以由blockDim.x * gridDim.x得到，现在N可以工作在任意值下了。