一、直观的感觉线程、线程块、线程格

为了直观的感觉线程、线程块、线程格，画了下面一个示意图。分为了两部分，一部分为线程格，另一部分为线程块，在图中线程格和线程块都画成了3维的，实际也可以是一维或者二维的。其中线程格里面最小的单元为线程块，而一个线程块里面最小的单元为线程。

二、threadIdx、blockIdx、blockDim和gridDim

可以把线程格和线程块都看作一个三维的矩阵。这里假设线程格是一个3*4*5的三维矩阵， 线程块是一个4*5*6的三维矩阵。

1. gridDim

   gridDim.x、gridDim.y、gridDim.z分别表示线程格各个维度的大小，所以有

   gridDim.x=3
   gridDim.y=4
   gridDim.z=5
   

2. blockDim
   blockDim.x、blockDim.y、blockDim.z分别表示线程块中各个维度的大小，所以有

   blockDim.x=4
   blockDim.y=5
   blockDim.z=6
   

3. blockIdx

   blockIdx.x、blockIdx.y、blockIdx.z分别表示当前线程块所处的线程格的坐标位置

4. threadIdx

   threadIdx.x、threadIdx.y、threadIdx.z分别表示当前线程所处的线程块的坐标位置

通过 blockIdx.x、blockIdx.y、blockIdx.z、threadIdx.x、threadIdx.y、threadIdx.z就可以完全定位一个线程的坐标位置了。

线程格里面总的线程个数N即可通过下面的公式算出

N = gridDim.x*gridDim.y*gridDim.z*blockDim.x*blockDim.y*blockDim.z

三、举例

将所有的线程排成一个序列，序列号为$0,1,2,\dots ,N$，如何找到当前的序列号？

1. 先找到当前线程位于线程格中的哪一个线程块blockId

   blockId = blockIdx.x + blockIdx.y*gridDim.x + blockIdx.z*gridDim.x*gridDim.y;
   

2. 找到当前线程位于线程块中的哪一个线程threadId

   threadId = threadIdx.x + threadIdx.y*blockDim.x + threadIdx.z*blockDim.x*blockDim.y;
   

3. 计算一个线程块中一共有多少个线程M

   M = blockDim.x*blockDim.y*blockDim.z
   

4. 求得当前的线程序列号idx

   idx = threadId + M*blockId;
   

下面是通过GPU并行计算实现的两个向量相减的例子

#include "cuda_runtime.h"
#include "device_launch_parameters.h"#include <stdio.h>
#include <stdlib.h>
#include <iostream>using namespace std;//block-thread 3D-3D
__global__ void testBlockThread9(int *c, const int *a, const int *b)
{int threadId_3D = threadIdx.x + threadIdx.y*blockDim.x + threadIdx.z*blockDim.x*blockDim.y;int blockId_3D = blockIdx.x + blockIdx.y*gridDim.x + blockIdx.z*gridDim.x*gridDim.y;int i = threadId_3D + (blockDim.x*blockDim.y*blockDim.z)*blockId_3D;c[i] = b[i] - a[i];
}void addWithCuda(int *c, const int *a, const int *b, unsigned int size)
{int *dev_a = 0;int *dev_b = 0;int *dev_c = 0;cudaSetDevice(0);cudaMalloc((void**)&dev_c, size * sizeof(int));cudaMalloc((void**)&dev_a, size * sizeof(int));cudaMalloc((void**)&dev_b, size * sizeof(int));cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);uint3 s1; s1.x = 5; s1.y = 2; s1.z = 2;uint3 s2; s2.x = size / 200; s2.y = 5; s2.z = 2;testBlockThread9<<<s1, s2 >>>(dev_c, dev_a, dev_b);cudaMemcpy(c, dev_c, size*sizeof(int), cudaMemcpyDeviceToHost);cudaFree(dev_a);cudaFree(dev_b);cudaFree(dev_c);cudaGetLastError();
}int main()
{const int n = 1000;int *a = new int[n];int *b = new int[n];int *c = new int[n];int *cc = new int[n];for (int i = 0; i < n; i++){a[i] = rand() % 100;b[i] = rand() % 100;c[i] = b[i] - a[i];}addWithCuda(cc, a, b, n);FILE *fp = fopen("out.txt", "w");for (int i = 0; i < n; i++)fprintf(fp, "%d %d\n", c[i], cc[i]);fclose(fp);bool flag = true;for (int i = 0; i < n; i++){if (c[i] != cc[i]){flag = false;break;}}if (flag == false)printf("no pass");elseprintf("pass");cudaDeviceReset();delete[] a;delete[] b;delete[] c;delete[] cc;getchar();return 0;
}