cuda中threadIdx、blockIdx、blockDim和gridDim的使用

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

为了直观的感觉线程、线程块、线程格，画了下面一个示意图。分为了两部分，一部分为线程格，另一部分为线程块，在图中线程格和线程块都画成了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");
    else
        printf("pass");

    cudaDeviceReset();

    delete[] a;
    delete[] b;
    delete[] c;
    delete[] cc;

    getchar();
    return 0;
}