CUDA线程索引的计算

版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。

本文链接： https://blog.csdn.net/hujingshuang/article/details/53097222

文章在如上博主分析下添加自己的理解：

#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <iostream>
 
using namespace std;
 
// 二：线程执行代码
__global__ void vector_add(float* vec1, float* vec2, float* vecres, int length) {
    int tid = threadIdx.x;
    if (tid < length) {
        vecres[tid] = vec1[tid] + vec2[tid];
    }
}
 
int main() {
    const int length = 16;                                      // 数组长度为16
    float a[length], b[length], c[length];                      // host中的数组
    for (int i = 0; i < length; i++) {                          // 初始赋值
        a[i] = b[i] = i;
    }
    float* a_device, *b_device, *c_device;                      // device中的数组
 
    cudaMalloc((void**)&a_device, length * sizeof(float));      // 分配内存
    cudaMalloc((void**)&b_device, length * sizeof(float));
    cudaMalloc((void**)&c_device, length * sizeof(float));
 
    cudaMemcpy(a_device, a, length * sizeof(float), cudaMemcpyHostToDevice);    // 将host数组的值拷贝给device数组
    cudaMemcpy(b_device, b, length * sizeof(float), cudaMemcpyHostToDevice);
 
    // 一：参数配置
    dim3 grid(1, 1, 1), block(length, 1, 1);                    // 设置参数
    vector_add<<<grid,block>>>(a_device, b_device, c_device, length);           // 启动kernel
 
    cudaMemcpy(c, c_device, length * sizeof(float), cudaMemcpyDeviceToHost);    // 将结果拷贝到host
 
    for (int i = 0; i < length; i++) {                          // 打印出来方便观察
        cout << c[i] << " ";
    }
    cout << endl;
 
    system("pause");
    return 0;
}

 

 从上面的代码实现中我们可以看出数组的大小是16，CUDA实现也使用了16个线程去计算。

dim3 grid(1, 1, 1), block(length, 1, 1);                    // 设置参数

先说grid，在这段代码中，我们设置参数为线程格（grid）中只有一个一维的block，该block的x维度上有16个，这个应该一下就看出来啦。因为grid(x,y,z)中的x=1,y=1,z=1，即各个维度均为1，所以是一维的，数量为x*y*z=1*1*1=1。如果没明白，再看两个例子：

dim3 grid1(2, 1, 1); // x=2, y=1, z=1
dim3 grid2(4, 2, 1); // x=4, y=2, z=1
dim3 grid3(2, 3, 4); // x=2, y=3, z=4

可以知道，grid1是一维的（因为y,z维度是1），grid2是二维的（因为z维度是1），grid3是三维的，且grid1,grid2,grid3中分别有2、8、24个block。同理，对于线程块（block），我们知道之前的代码中，block中存在16个线程，且该线程块维度是一维的，因为block(x,y,z)中x=length=16,y=1,z=1。
我画个图来帮助理解，大概就是这样子的：

dim3 grid(1, 1, 1), block(length, 1, 1);                    // 设置参数

dim3 grid(1, 1, 1), block(length, 1, 1);                    // 设置参数

 

   OK，我想这下应该就清楚了，就是一个一维的block（此处只有x维度上存在16个线程）。所以，內建变量只有一个在起作用，就是threadIdx.x，它的范围是[0,15]。因此，我们在计算线程索引是，只用这个內建变量就行了（其他的为0，写了也不起作用）：

// 二：线程执行代码
__global__ void vector_add(float* vec1, float* vec2, float* vecres, int length) {
    int tid = threadIdx.x;              // 只使用了threadIdx.x，假设threadIdx.x=0,那么使用0号线程；threadIdx.x=1,使用1号线程；threadIdx.x=15，使用15号线程（物理上是第16个线程，因为线程计算从0开始）
    if (tid < length) {
        vecres[tid] = vec1[tid] + vec2[tid];
    }
}

OK，看到这里，你可能还是不大明白什么一维二维的，我们再来看一个：

dim3 grid(1, 1, 1), block(8, 2, 1);                    // 设置参数

 

 根据上面的介绍，我们知道这个线程格只有一个一维的线程块，该线程块内的线程是二维的，x的维度为8，y的维度为2，共有8*2=16个线程，如果要用这16个线程来计算数组的累加，当然是可以的，但是我们这里需要改动一下线程执行代码中的索引计算方式了。

// 二：线程执行代码
__global__ void vector_add(float* vec1, float* vec2, float* vecres, int length) {
    int tid = threadIdx.y * blockDim.x +  threadIdx.x;  // 使用了threadIdx.x, threadIdx.x, blockDim.x；
公式计算可以理解为：threadIdx.y=0，threadIdx.x=0.计算出tid=0，使用了0号线程；threadIdx.y=0，threadIdx.x=1，tid=1，使用了1号线程（这些都是第0个线程块中的）；
每层跳跃的个数blockDim.x=8，如果threadIdx.y=1,threadIdx.x=0，计算出1*8+0=8，其实是跳到了下一个             
线程块的开始的第一个线程（但是物理上没有这么划分的，这只是逻辑上的划分，物理上就是第8号线程也就是第9个线程，因为线程号是从0开始计算的）；
再看一个如果threadIdx.y=1（说明已经在1号thread块中了），threadIdx.x=3，tid=1*8+3=11，对应到第一个thread块中的3号线程，似乎加多少就是几号
线程（偏移量），那么可以理解为threadIdx.y是第几个线程块（block），threadIdx.x是几个线程块中的线程。

    if (tid < length) {
        vecres[tid] = vec1[tid] + vec2[tid];
    }
}

我们一定要有并行思想，这里有16个线程，kernel启动后，每个线程都有自己的索引号，比如某个线程位于grid中哪个维度的block（即blockIdx.x,block.y,block.z），又位于该block的哪个维度的线程（即threadIdx.x,threadIdx.y,threadIdx.z），利用这些线程索引号映射到对应的数组下标，我们要做的工作就是将保证这些下标不重复（如果重复的话，那就惨了），最初那种一维的计算方式就不行了。因此，通过使用threadIdx，blockDim来进行映射（偏移）。blockDim.x=8,blockDim.y=2，如上代码。
  其实，我感觉有些我不能用文字准确、清晰的描述出来，所以咯，我们再来一个例子吧，我相信，多看一看，多想一想就明白了。

dim3 grid(1, 1, 1), block(4, 4, 1);                    // 设置参数

 我们将block改成上面的这样，其线程模型为下图：

 

 当然，kernel函数的代码依然可以不用变动，这个应该想得清楚，还是再写一下吧。

// 二：线程执行代码
__global__ void vector_add(float* vec1, float* vec2, float* vecres, int length) {
    int tid = threadIdx.y * blockDim.x +  threadIdx.x;  // 使用了threadIdx.x, threadIdx.x, blockDim.x；
继续举例说明：如果threadIdx.y=3,threadIdx.x=2,那么tid=3*4+2=14，物理上是第14号线程（14号线程是第15个）（因为线程计算方式从0号开始的所以16个线程最多表示到15），逻辑上是第3个线程
块的第2号线程，依次类比理解即可。

    if (tid < length) {
        vecres[tid] = vec1[tid] + vec2[tid];
    }
}

以上内容我们分别介绍了用一维和二维线程来计算一维数组的求和，实际上数组的维度与线程格、线程块和线程的维度并不是那么密不可分的，都可以组合实现，只不过在实现时，良好的参数配置对索引的计算很方便，而且由于grid、block、thread维度的限制，还有warpSize的限制，所以对于较大的数据量来说，我们应该做到心中有数，进行有效的块分解。

 

 现在来看看二维的block，在整个文章中，我只讲解一维、二维的，因为三维的我不知道怎么画图啦，而且不好描述，免得误导大家。
 还是上面的一维数组，长度为16。

dim3 grid(16, 1, 1), block(1, 1, 1);                    // 设置参数

先来个线程模型图，我想大家并不会感到惊讶，绿色的区域表示grid，蓝色的区域表示block，图中有一个grid和16个block，每个block都是一维，而且x维度上只有一个线程的：

 

 显然，我们的线程索引代码应该为如下：

// 二：线程执行代码
__global__ void vector_add(float* vec1, float* vec2, float* vecres, int length) {
    int tid = blockIdx.x;
    if (tid < length) {
        vecres[tid] = vec1[tid] + vec2[tid];
    }
}

或许你会有疑惑，那么我们再来看一个：

dim3 grid(4, 1, 1), block(4, 1, 1);

 

 

 线程索引代码应该为如下：

// 二：线程执行代码
__global__ void vector_add(float* vec1, float* vec2, float* vecres, int length) {
    int tid = blockIdx.x * gridDim.x + threadIdx.x;//一维的是threadIdx.y*blockDim.x，二维的变成blockIdx.x*gridDim.x，维度整体上升一个；举例论证：blockIdx.x=2，threadIdx.x=2,tid=2*4+2=10
(第10号线程，物理上是第11个线程，逻辑上是第2个线程块的第2个线程)
    if (tid < length) {
        vecres[tid] = vec1[tid] + vec2[tid];
    }
}

再比如：

dim3 grid(2, 2, 1), block(2, 2, 1);

 

 那么执行代码及索引计算如下：

// 二：线程执行代码
__global__ void vector_add(float* vec1, float* vec2, float* vecres, int length) {
    // 在第几个块中 * 块的大小 + 块中的x, y维度（几行几列），（blockDim.x*blockDim.y相当于前面的gridDim.x相当于前面的blockDim.x等同于跳过的快的大小），然后从高维加到低维。举例论证：blockIdx.y=1,blockIdx.x=1,那么
第一个括号里面（1*2+1=3），第二个括号（2*2=4），threadIdx.y=1,threadIdx.x=1,加号后面等于1*2+1=3，所以tid=3*4+3=15（第15号线程，物理上是第16个，逻辑上是第（1，1）线程块里面的（1，1）个线程。）

    int tid = (blockIdx.y * gridDim.x + blockIdx.x) * (blockDim.x * blockDim.y) + threadIdx.y * blockDim.x + threadIdx.x;
    if (tid < length) {
        vecres[tid] = vec1[tid] + vec2[tid];
    }
}

最后再来一发，我给个图，我们来倒推其参数及相关执行代码，如下：

 

 

dim3 grid(8, 4, 1), block(8, 2, 1);

共有8*4*8*2=512个线程，当然在CUDA编程中，这算很少的了。如果是一幅512x512大小的图像做加或点乘之类的运算，随随便便就是几十万的线程数了。
万变不离其宗，第一个是一维的计算方式，第二个是二维的计算公式：

__global__ void vector_add(float* vec1, float* vec2, float* vecres, int length) {
    // 在第几个块中 * 块的大小 + 块中的x, y维度（几行几列）
    int tid = (blockIdx.y * gridDim.x + blockIdx.x) * (blockDim.x * blockDim.y) + threadIdx.y * blockDim.y + threadIdx.x;
    if (tid < length) {
        vecres[tid] = vec1[tid] + vec2[tid];
    }
}

__global__ void vector_add(float** mat1, float** mat2, float** matres, int width) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    if (x < width && y < width) {
        matres[x][y] = mat1[x][y] + mat2[x][y];
    }
}