今天看了《GPU高性能编程》中的矢量求和的例子，对其中的并行线程不是很理解，源码为：
//a+b
#include"…/common/book.h"
#define N 331024
global void add(int a,int b,int c)
{
int tid=threadIdx.x+blockIdx.x * blockDim.x;
while(tid<N)
{
c[tid]=a[tid]+b[tid];
tid+=blockDim.x * gridDim.x;
}
}
int main(void)
{
int a[N],b[N],c[N];
int dev_a,dev_b,dev_c;
//在GPU上分配内存
cudaMalloc((void)&dev_a,N * sizeof(int));
cudaMalloc((void)&dev_b,N * sizeof(int));
cudaMalloc((void)&dev_c,N * sizeof(int));
//在CPU上为数组a、b赋值
for(int i=0;i<N;i++)
{
a[i]=i;
b[i]=ii;
}
//将数组a、b复制到GPU
cudaMemcpy(dev_a,a,N * sizeof(int),cudaMemcpyHostToDevice);
cudaMemcpy(dev_b,b,N * sizeof(int),cudaMemcpyHostToDevice);
add<<<128,128>>>(dev_a,dev_b,dev_c);
//将数组c从GPU复制到CPU
cudaMemcpy(c,dev_c,N * sizeof(int),cudaMemcpyDeviceToHost);
bool success=true;
for(int i=0;i<N;i++)
{
if((a[i]+b[i])!=c[i])
{
printf(“Error:%d + %d != %d\n”,a[i],b[i],c[i]);
success=false;
}
}
if(success)
printf(“We did it!\n”);
cudaFree(dev_a);
cudaFree(dev_b);
cudaFree(dev_c);
return 0;
}
其中 tid+=blockDim.x * gridDim.x;为什么每次循环tid都要加上blockDim.x * gridDim.x。
因为你有33K组数据需要相加。
但是你却一共只有128 * 128个线程(即：16K个).
所以这些线程都算完前16K个后，又集体跃迁了2次。来完成后续的17K次。
请问blockDim.x * gridDim.x的值是多少，blockDim.x * gridDim.x是代表当前运行的线程数，还是代表线程格中线程的总数呢
请LZ查阅programming guide的Appendix B里面关于内建数据类型的介绍。
blockDim.x和gridDim.x的含义也请查阅上述章节。
您的示例代码中，这两个值都是128。至于为何是128，请您搞清楚调用kernel时参数的含义，此问题不问自明。
在您给出的例子中，表示总的线程数。
但是，请您明白，这两个数值的乘积，在其他形状的grid/block划分下，未必具有这个含义。
希望您能明白，祝您编码愉快~
（1）blockDim.x和gridDim.x都是128，请仔细看你的代码。（“<<<128,128>>>”这个地方）。
（2）你理解的正确。这2个数字，是(grid里的)block和(block里的)thread的总数。而不是当前存活的。
如果依然不理解，我推荐您参考如下文档：
《 CUDA C/C++ Programming Guide》的第二章。（也就是越过前面介绍部分，即可看到）。
是不是可以这样理解 总共要计算的数据量为331024，但所给的一个grid总共的线程数为128128，先从1到128*128个线程先运行一次，然后每个线程在循环的时候加上blockDim.x * gridDim.x，得到一个新的线程号，也就是本grid中的线程已用完，分配了另外一个grid中的线程，这样理解对吗？
这样理解是错误的，我给您用个比喻来说明下：
你有33块砖头(用来比喻你的33K组数据), 以及你有16个手下(用来比喻你的grid里的16K个线程)，
显然这16个手下决定这么搬运：
（1）他们先搬运第1-16块。（这16个手下，每人对付的砖头编号分别为1到16）
（2）然后搬运第17-32块。（这16个手下，每人搬运自己原来编号加上16的那块砖头，实现搬运第17到32砖头）
（3）然后搬运第33块。（每个手下搬运的砖头编号再度+16，并只搬运存在的砖头。不存在的无视（你的if))
这样好理解吗？
您的理解是不正确的
1：您混淆了线程号和寻址。是的，线程号可以辅助寻址，而且我们一般也是这么干的，但是，代码中每次循环加入的偏移量只是为了寻址，而不是“本grid中的线程已用完，分配了另外一个grid中的线程”。重新启动另外一个grid的线程，需要重新启动一个kernel，这一般只能在host端完成（SM 3.5之前的设备）或者在SM 3.5的设备上通过kernel invoke kernel来实现。
2：因为前面说过，循环中每次增加的只是偏移量，所以只需要维护每个线程找对自己的位置即可。和“究竟是所有线程都存活着，先都把第一次累加实现了，再去加第二次”，还是“只有部分线程存活，先把自己的活干完了，然后死掉，其他线程再上”没有关系，也不影响结果。一般的CUDA程序，线程数量会非常庞大，所以往往是后一种情况。
以上说法供您参考，祝您编码顺利~
system
14
嗯嗯，恭喜LZ EXP+5！
欢迎经常来论坛讨论问题~
祝您新年愉快~