CUDA线程层次一文搞懂｜参加CUDA线上训练营

设备术语

• Host：CPU 和 内存 (host memory)
• Device：GPU 和显存 (device memory)

CUDA 线程层次

CUDA 线程层次分为：

• Thread
  • 所有线程执行相同的核函数
  • 并行执行
• Thread Block
  • 执行在一个 Streaming Multiprocessor （SM）
  • 同一个 Block 中的线程可以协作
• Thread Grid
  • 一个 Grid当中的 Block 可以在多个 SM 中执行

CUDA执行顺序

1. 加载核函数
2. 将 Grid 分配到一个 Device
3. 根据 <<<..>>> 内的执行设置的第一个参数，Giga threads engine 将 block 分配到 SM 中。一个 Block 内的线程一定会在同一个 SM 内，一个 SM 可以有很多个 Block
4. 根据 <<<..>>> 内的执行设置的第二个参数，Warp 调度器会将调用线程
5. Warp 调度器为了提高运行效率，会将每 32 个线程分为一组，称作-个 warp
6. 每个 warp 会被分配到 32 个 core 上运行

CUDA 的一切精髓就是并行加速冲冲冲！

如何计算索引

首先来看看基本概念：

• threadIdx.[x y z]是执行当前kernel函数的线程在block中的索引值（threadIdx.x是1，threadIdx.y是0）

• blockIdx.[x y z]是指执行当前kernel函数的线程所在block，在grid中的索引值（blockIdx.x是1，blockIdx.y是1）

• blockDim.[x y z]表示一个block中包含多少个线程（blockDim.x是5，blockDim.y是3）

• gridDim.[x y z]表示一个grid中包含多少个block（gridDim.x是3，gridDim.y是2）

计算矩阵运算的时候，将矩阵中的一行取出来，但是因为 CUDA 是多个线程并行的，就是每个线程里面都会同时获取到矩阵行中的某个元素，我们就需要在核函数里面计算出来这个元素在原来矩阵行中的索引，下面是个例子：

Demo

接下来，我们通过完成一个向量加法的实例来实践一下: 。

为了完成这个程序，我们先要将数据传输给GPU，并在GPU完成计算的时候，将数据从GPU中传输给CPU内存。这时我们就需要考虑如何申请GPU存储单元，以及内存和显存之前的数据传输。

我们利用cudaMalloc()来进行GPU存储单元的申请，利用cudaMemcpy()来完成数据的传输

代码如下：

#include <math.h>
#include <stdio.h>void __global__ add(const double *x, const double *y, double *z, int count)
{const int n = blockDim.x * blockIdx.x + threadIdx.x;if( n < count){z[n] = x[n] + y[n];}}
void check(const double *z, const int N)
{bool error = false;for (int n = 0; n < N; ++n){if (fabs(z[n] - 3) > (1.0e-10)){error = true;}}printf("%s\n", error ? "Errors" : "Pass");
}int main(void)
{const int N = 1000;const int M = sizeof(double) * N;double *h_x = (double*) malloc(M);double *h_y = (double*) malloc(M);double *h_z = (double*) malloc(M);for (int n = 0; n < N; ++n){h_x[n] = 1;h_y[n] = 2;}double *d_x, *d_y, *d_z;cudaMalloc((void **)&d_x, M);cudaMalloc((void **)&d_y, M);cudaMalloc((void **)&d_z, M);cudaMemcpy(d_x, h_x, M, cudaMemcpyHostToDevice);cudaMemcpy(d_y, h_y, M, cudaMemcpyHostToDevice);const int block_size = 128;const int grid_size = (N + block_size - 1) / block_size;add<<<grid_size, block_size>>>(d_x, d_y, d_z, N);cudaMemcpy(h_z, d_z, M, cudaMemcpyDeviceToHost);check(h_z, N);free(h_x);free(h_y);free(h_z);cudaFree(d_x);cudaFree(d_y);cudaFree(d_z);return 0;
}